[clang] [OpenACC] implement loop 'worker' clause. (PR #112206)

Erich Keane via cfe-commits cfe-commits at lists.llvm.org
Mon Oct 14 07:07:07 PDT 2024


https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/112206

The worker clause specifies iterations of the loop/ that are executed in parallel by distributing the iterations among the multiple works within a single gang.

The sema rules for this type are simply that it cannot be combined with a `kernel` construct with a `num_workers` clause, child `loop` clauses cannot contain a `gang` or `worker` clause, and that the argument is oly allowed when associated with a `kernel`.

>From bc9b06006b15a4f7f3bd4b89d9a5191a43280e04 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Fri, 11 Oct 2024 09:59:24 -0700
Subject: [PATCH] [OpenACC] implement loop 'worker' clause.

The worker clause specifies iterations of the loop/ that are executed in
parallel by distributing the iterations among the multiple works within
a single gang.

The sema rules for this type are simply that it cannot be combined with
a `kerne` construct with a `num_workers` clause, child `loop` clauses
cannot contain a `gang` or `worker` clause, and that the argument is oly
allowed when associated with a `kernel`.
---
 clang/include/clang/AST/OpenACCClause.h       |  42 ++-
 .../clang/Basic/DiagnosticSemaKinds.td        |  21 +-
 clang/include/clang/Basic/OpenACCClauses.def  |   1 +
 clang/include/clang/Sema/SemaOpenACC.h        |  16 ++
 clang/lib/AST/OpenACCClause.cpp               |  30 +-
 clang/lib/AST/StmtProfile.cpp                 |   6 +
 clang/lib/AST/TextNodeDumper.cpp              |   1 +
 clang/lib/Parse/ParseOpenACC.cpp              |   1 +
 clang/lib/Sema/SemaOpenACC.cpp                | 125 +++++++-
 clang/lib/Sema/TreeTransform.h                |  28 ++
 clang/lib/Serialization/ASTReader.cpp         |   7 +-
 clang/lib/Serialization/ASTWriter.cpp         |   9 +-
 .../AST/ast-print-openacc-loop-construct.cpp  |  39 +++
 clang/test/ParserOpenACC/parse-clauses.c      |  19 +-
 .../compute-construct-device_type-clause.c    |   3 +-
 ...p-construct-auto_seq_independent-clauses.c |  15 +-
 .../loop-construct-device_type-clause.c       |   1 -
 .../SemaOpenACC/loop-construct-worker-ast.cpp | 270 ++++++++++++++++++
 .../loop-construct-worker-clause.cpp          | 202 +++++++++++++
 clang/tools/libclang/CIndex.cpp               |   5 +
 20 files changed, 766 insertions(+), 75 deletions(-)
 create mode 100644 clang/test/SemaOpenACC/loop-construct-worker-ast.cpp
 create mode 100644 clang/test/SemaOpenACC/loop-construct-worker-clause.cpp

diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index f3a09eb651458d..e8b8f477f91ae7 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -145,32 +145,6 @@ class OpenACCVectorClause : public OpenACCClause {
   }
 };
 
-// Not yet implemented, but the type name is necessary for 'seq' diagnostics, so
-// this provides a basic, do-nothing implementation. We still need to add this
-// type to the visitors/etc, as well as get it to take its proper arguments.
-class OpenACCWorkerClause : public OpenACCClause {
-protected:
-  OpenACCWorkerClause(SourceLocation BeginLoc, SourceLocation EndLoc)
-      : OpenACCClause(OpenACCClauseKind::Worker, BeginLoc, EndLoc) {
-    llvm_unreachable("Not yet implemented");
-  }
-
-public:
-  static bool classof(const OpenACCClause *C) {
-    return C->getClauseKind() == OpenACCClauseKind::Worker;
-  }
-
-  static OpenACCWorkerClause *
-  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 a clause that has a list of parameters.
 class OpenACCClauseWithParams : public OpenACCClause {
   /// Location of the '('.
@@ -541,6 +515,22 @@ class OpenACCGangClause final
          ArrayRef<Expr *> IntExprs, SourceLocation EndLoc);
 };
 
+class OpenACCWorkerClause : public OpenACCClauseWithSingleIntExpr {
+protected:
+  OpenACCWorkerClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                      Expr *IntExpr, SourceLocation EndLoc);
+
+public:
+  static bool classof(const OpenACCClause *C) {
+    return C->getClauseKind() == OpenACCClauseKind::Worker;
+  }
+
+  static OpenACCWorkerClause *Create(const ASTContext &Ctx,
+                                     SourceLocation BeginLoc,
+                                     SourceLocation LParenLoc, Expr *IntExpr,
+                                     SourceLocation EndLoc);
+};
+
 class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr {
   OpenACCNumWorkersClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
                           Expr *IntExpr, SourceLocation EndLoc);
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 3c62a017005e59..00201a7f192f17 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12686,21 +12686,22 @@ def err_acc_intervening_code
 def err_acc_gang_multiple_elt
     : Error<"OpenACC 'gang' clause may have at most one %select{unnamed or "
             "'num'|'dim'|'static'}0 argument">;
-def err_acc_gang_arg_invalid
-    : Error<"'%0' argument on 'gang' clause is not permitted on a%select{n "
-            "orphaned|||}1 'loop' construct %select{|associated with a "
+def err_acc_int_arg_invalid
+    : Error<"'%1' argument on '%0' clause is not permitted on a%select{n "
+            "orphaned|||}2 'loop' construct %select{|associated with a "
             "'parallel' compute construct|associated with a 'kernels' compute "
-            "construct|associated with a 'serial' compute construct}1">;
+            "construct|associated with a 'serial' compute construct}2">;
 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_gang_num_gangs_conflict
-    : Error<"'num' argument to 'gang' clause not allowed on a 'loop' construct "
-            "associated with a 'kernels' construct that has a 'num_gangs' "
+def err_acc_num_arg_conflict
+    : Error<"'num' argument to '%0' clause not allowed on a 'loop' construct "
+            "associated with a 'kernels' construct that has a "
+            "'%select{num_gangs|num_workers}1' "
             "clause">;
-def err_acc_gang_inside_gang
-    : Error<"loop with a 'gang' clause may not exist in the region of a 'gang' "
-            "clause on a 'kernels' compute construct">;
+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 'kernels' compute construct}2">;
 
 // AMDGCN builtins diagnostics
 def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 2a098de31eb618..4c0b56dc13e625 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -56,6 +56,7 @@ VISIT_CLAUSE(Seq)
 VISIT_CLAUSE(Tile)
 VISIT_CLAUSE(VectorLength)
 VISIT_CLAUSE(Wait)
+VISIT_CLAUSE(Worker)
 
 #undef VISIT_CLAUSE
 #undef CLAUSE_ALIAS
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 59a9648d5f9380..e253610a84b0bf 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -118,6 +118,11 @@ class SemaOpenACC : public SemaBase {
   /// 'kernel' construct, this will have the source location for it. This
   /// permits us to implement the restriction of no further 'gang' clauses.
   SourceLocation LoopGangClauseOnKernelLoc;
+  /// If there is a current 'active' loop construct with a 'worker' clause on it
+  /// (on any sort of construct), this has the source location for it.  This
+  /// permits us to implement the restriction of no further 'gang' or 'worker'
+  /// clauses.
+  SourceLocation LoopWorkerClauseLoc;
 
   // Redeclaration of the version in OpenACCClause.h.
   using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
@@ -224,11 +229,15 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::Tile ||
+              ClauseKind == OpenACCClauseKind::Worker ||
+              ClauseKind == OpenACCClauseKind::Vector ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
 
       // 'async' and 'wait' have an optional IntExpr, so be tolerant of that.
       if ((ClauseKind == OpenACCClauseKind::Async ||
+           ClauseKind == OpenACCClauseKind::Worker ||
+           ClauseKind == OpenACCClauseKind::Vector ||
            ClauseKind == OpenACCClauseKind::Wait) &&
           std::holds_alternative<std::monostate>(Details))
         return 0;
@@ -271,6 +280,8 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::Tile ||
               ClauseKind == OpenACCClauseKind::Gang ||
+              ClauseKind == OpenACCClauseKind::Worker ||
+              ClauseKind == OpenACCClauseKind::Vector ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
 
@@ -401,6 +412,8 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::Tile ||
+              ClauseKind == OpenACCClauseKind::Worker ||
+              ClauseKind == OpenACCClauseKind::Vector ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
       Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}};
@@ -410,6 +423,8 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::Tile ||
+              ClauseKind == OpenACCClauseKind::Worker ||
+              ClauseKind == OpenACCClauseKind::Vector ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
       Details = IntExprDetails{std::move(IntExprs)};
@@ -663,6 +678,7 @@ class SemaOpenACC : public SemaBase {
     ComputeConstructInfo OldActiveComputeConstructInfo;
     OpenACCDirectiveKind DirKind;
     SourceLocation OldLoopGangClauseOnKernelLoc;
+    SourceLocation OldLoopWorkerClauseLoc;
     llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
     LoopInConstructRAII LoopRAII;
 
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 6fb8fe0b8cfeef..638252fd811f1d 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -44,7 +44,8 @@ bool OpenACCClauseWithCondition::classof(const OpenACCClause *C) {
 bool OpenACCClauseWithSingleIntExpr::classof(const OpenACCClause *C) {
   return OpenACCNumWorkersClause::classof(C) ||
          OpenACCVectorLengthClause::classof(C) ||
-         OpenACCCollapseClause::classof(C) || OpenACCAsyncClause::classof(C);
+         OpenACCWorkerClause::classof(C) || OpenACCCollapseClause::classof(C) ||
+         OpenACCAsyncClause::classof(C);
 }
 OpenACCDefaultClause *OpenACCDefaultClause::Create(const ASTContext &C,
                                                    OpenACCDefaultClauseKind K,
@@ -403,11 +404,24 @@ OpenACCGangClause::Create(const ASTContext &C, SourceLocation BeginLoc,
       OpenACCGangClause(BeginLoc, LParenLoc, GangKinds, IntExprs, EndLoc);
 }
 
+OpenACCWorkerClause::OpenACCWorkerClause(SourceLocation BeginLoc,
+                                         SourceLocation LParenLoc,
+                                         Expr *IntExpr, SourceLocation EndLoc)
+    : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Worker, BeginLoc,
+                                     LParenLoc, IntExpr, EndLoc) {
+  assert((!IntExpr || IntExpr->isInstantiationDependent() ||
+          IntExpr->getType()->isIntegerType()) &&
+         "Int expression type not scalar/dependent");
+}
+
 OpenACCWorkerClause *OpenACCWorkerClause::Create(const ASTContext &C,
                                                  SourceLocation BeginLoc,
+                                                 SourceLocation LParenLoc,
+                                                 Expr *IntExpr,
                                                  SourceLocation EndLoc) {
-  void *Mem = C.Allocate(sizeof(OpenACCWorkerClause));
-  return new (Mem) OpenACCWorkerClause(BeginLoc, EndLoc);
+  void *Mem =
+      C.Allocate(sizeof(OpenACCWorkerClause), alignof(OpenACCWorkerClause));
+  return new (Mem) OpenACCWorkerClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
 }
 
 OpenACCVectorClause *OpenACCVectorClause::Create(const ASTContext &C,
@@ -638,3 +652,13 @@ void OpenACCClausePrinter::VisitGangClause(const OpenACCGangClause &C) {
     OS << ")";
   }
 }
+
+void OpenACCClausePrinter::VisitWorkerClause(const OpenACCWorkerClause &C) {
+  OS << "worker";
+
+  if (C.hasIntExpr()) {
+    OS << "(num: ";
+    printExpr(C.getIntExpr());
+    OS << ")";
+  }
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 6161b1403ed35d..25b1cbb8590869 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2629,6 +2629,12 @@ void OpenACCClauseProfiler::VisitAsyncClause(const OpenACCAsyncClause &Clause) {
     Profiler.VisitStmt(Clause.getIntExpr());
 }
 
+void OpenACCClauseProfiler::VisitWorkerClause(
+    const OpenACCWorkerClause &Clause) {
+  if (Clause.hasIntExpr())
+    Profiler.VisitStmt(Clause.getIntExpr());
+}
+
 void OpenACCClauseProfiler::VisitWaitClause(const OpenACCWaitClause &Clause) {
   if (Clause.hasDevNumExpr())
     Profiler.VisitStmt(Clause.getDevNumExpr());
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index ac8c196777f9b8..beccb0615f0e9c 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -420,6 +420,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
     case OpenACCClauseKind::Self:
     case OpenACCClauseKind::Seq:
     case OpenACCClauseKind::Tile:
+    case OpenACCClauseKind::Worker:
     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.
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 635039b724e6a0..51d4dc38c17f67 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -1130,6 +1130,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
           Parens.skipToEnd();
           return OpenACCCanContinue();
         }
+        ParsedClause.setIntExprDetails(IntExpr.get());
         break;
       }
       case OpenACCClauseKind::Async: {
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 30d73d621db69b..1b24331cbd87ca 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -377,6 +377,18 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
     default:
       return false;
     }
+  case OpenACCClauseKind::Worker: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Loop:
+    case OpenACCDirectiveKind::ParallelLoop:
+    case OpenACCDirectiveKind::SerialLoop:
+    case OpenACCDirectiveKind::KernelsLoop:
+    case OpenACCDirectiveKind::Routine:
+      return true;
+    default:
+      return false;
+    }
+  }
   }
 
   default:
@@ -500,7 +512,6 @@ class SemaOpenACCClauseVisitor {
 
   OpenACCClause *Visit(SemaOpenACC::OpenACCParsedClause &Clause) {
     switch (Clause.getClauseKind()) {
-    case OpenACCClauseKind::Worker:
     case OpenACCClauseKind::Vector: {
       // TODO OpenACC: These are only implemented enough for the 'seq'
       // diagnostic, otherwise treats itself as unimplemented.  When we
@@ -1024,6 +1035,75 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause(
                                           Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
+    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)
+    return isNotImplemented();
+
+  Expr *IntExpr =
+      Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr;
+
+  if (IntExpr) {
+    switch (SemaRef.getActiveComputeConstructInfo().Kind) {
+    case OpenACCDirectiveKind::Invalid:
+      SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
+          << OpenACCClauseKind::Worker << "num" << /*orphan=*/0;
+      IntExpr = nullptr;
+      break;
+    case OpenACCDirectiveKind::Parallel:
+      SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
+          << OpenACCClauseKind::Worker << "num" << /*parallel=*/1;
+      IntExpr = nullptr;
+      break;
+    case OpenACCDirectiveKind::Serial:
+      SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
+          << OpenACCClauseKind::Worker << "num" << /*serial=*/3;
+      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 << /*num_workers=*/1;
+        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;
+  }
+
+  return OpenACCWorkerClause::Create(Ctx, Clause.getBeginLoc(),
+                                     Clause.getLParenLoc(), IntExpr,
+                                     Clause.getEndLoc());
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   if (DiagIfSeqClause(Clause))
@@ -1061,8 +1141,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
                         llvm::IsaPred<OpenACCNumGangsClause>);
 
       if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
-        SemaRef.Diag(ER.get()->getBeginLoc(),
-                     diag::err_acc_gang_num_gangs_conflict);
+        SemaRef.Diag(ER.get()->getBeginLoc(), diag::err_acc_num_arg_conflict)
+            << OpenACCClauseKind::Gang << /*num_gangs=*/0;
         SemaRef.Diag((*Itr)->getBeginLoc(),
                      diag::note_acc_previous_clause_here);
         continue;
@@ -1091,12 +1171,28 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
   if (SemaRef.LoopGangClauseOnKernelLoc.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_gang_inside_gang);
+    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+        << OpenACCClauseKind::Gang << OpenACCClauseKind::Gang
+        << /*kernels construct info*/ 1;
     SemaRef.Diag(SemaRef.LoopGangClauseOnKernelLoc,
                  diag::note_acc_previous_clause_here);
     return nullptr;
   }
 
+  // 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::Gang << OpenACCClauseKind::Worker
+        << /*kernels construct info*/ 1;
+    SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
+                 diag::note_acc_previous_clause_here);
+    return nullptr;
+  }
+
   return OpenACCGangClause::Create(Ctx, Clause.getBeginLoc(),
                                    Clause.getLParenLoc(), GangKinds, IntExprs,
                                    Clause.getEndLoc());
@@ -1216,6 +1312,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     ArrayRef<OpenACCClause *> Clauses)
     : SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo),
       DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc),
+      OldLoopWorkerClauseLoc(S.LoopWorkerClauseLoc),
       LoopRAII(SemaRef, /*PreserveDepth=*/false) {
   // Compute constructs end up taking their 'loop'.
   if (DirKind == OpenACCDirectiveKind::Parallel ||
@@ -1232,6 +1329,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     //
     // Implement the 'unless within a nested compute region' part.
     SemaRef.LoopGangClauseOnKernelLoc = {};
+    SemaRef.LoopWorkerClauseLoc = {};
   } else if (DirKind == OpenACCDirectiveKind::Loop) {
     SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
     SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
@@ -1252,6 +1350,12 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
       if (Itr != Clauses.end())
         SemaRef.LoopGangClauseOnKernelLoc = (*Itr)->getBeginLoc();
     }
+
+    if (UnInstClauses.empty()) {
+      auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCWorkerClause>);
+      if (Itr != Clauses.end())
+        SemaRef.LoopWorkerClauseLoc = (*Itr)->getBeginLoc();
+    }
   }
 }
 
@@ -1324,6 +1428,7 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
 SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
   SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
   SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc;
+  SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc;
 
   if (DirKind == OpenACCDirectiveKind::Parallel ||
       DirKind == OpenACCDirectiveKind::Serial ||
@@ -1934,8 +2039,8 @@ ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK, Expr *E) {
       // construct, or an orphaned loop construct, the gang clause behaves as
       // follows. ... The num argument is not allowed.
     case OpenACCGangKind::Num:
-      Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid)
-          << GK
+      Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
+          << OpenACCClauseKind::Gang << GK
           << (/*orphan/parallel=*/ActiveComputeConstructInfo.Kind ==
                       OpenACCDirectiveKind::Parallel
                   ? 1
@@ -1951,8 +2056,8 @@ ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK, Expr *E) {
     // construct, the gang clause behaves as follows. ... The dim argument is
     // not allowed.
     case OpenACCGangKind::Dim:
-      Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid)
-          << GK << /*kernels=*/2;
+      Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
+          << OpenACCClauseKind::Gang << GK << /*kernels=*/2;
       return ExprError();
     // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
     // construct, the gang clause behaves as follows. ... An argument with no
@@ -1975,8 +2080,8 @@ ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK, Expr *E) {
     // too, so we disallow them too.
     case OpenACCGangKind::Dim:
     case OpenACCGangKind::Num:
-      Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid)
-          << GK << /*Kernels=*/3;
+      Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
+          << OpenACCClauseKind::Gang << GK << /*Kernels=*/3;
       return ExprError();
     case OpenACCGangKind::Static:
       return CheckGangStaticExpr(*this, E);
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index cde40773336866..45e8b3cf6bd8fc 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11787,6 +11787,34 @@ void OpenACCClauseTransform<Derived>::VisitAsyncClause(
                                          : nullptr,
       ParsedClause.getEndLoc());
 }
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitWorkerClause(
+    const OpenACCWorkerClause &C) {
+  if (C.hasIntExpr()) {
+    // restrictions on this expression are all "does it exist in certain
+    // situations" that are not possible to be dependent, so the only check we
+    // have is that it transforms, and is an int expression.
+    ExprResult Res = Self.TransformExpr(const_cast<Expr *>(C.getIntExpr()));
+    if (!Res.isUsable())
+      return;
+
+    Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid,
+                                                C.getClauseKind(),
+                                                C.getBeginLoc(), Res.get());
+    if (!Res.isUsable())
+      return;
+    ParsedClause.setIntExprDetails(Res.get());
+  }
+
+  NewClause = OpenACCWorkerClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(),
+      ParsedClause.getNumIntExprs() != 0 ? ParsedClause.getIntExprs()[0]
+                                         : nullptr,
+      ParsedClause.getEndLoc());
+}
+
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitWaitClause(
     const OpenACCWaitClause &C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 31881a39de47f7..ecc5d3c59a3549 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12339,10 +12339,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCGangClause::Create(getContext(), BeginLoc, LParenLoc,
                                      GangKinds, Exprs, EndLoc);
   }
+  case OpenACCClauseKind::Worker: {
+    SourceLocation LParenLoc = readSourceLocation();
+    Expr *WorkerExpr = readBool() ? readSubExpr() : nullptr;
+    return OpenACCWorkerClause::Create(getContext(), BeginLoc, LParenLoc,
+                                       WorkerExpr, EndLoc);
+  }
 
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
-  case OpenACCClauseKind::Worker:
   case OpenACCClauseKind::Vector:
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 583d9a4bccb800..0a6e260e3e4e93 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8192,10 +8192,17 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     }
     return;
   }
+  case OpenACCClauseKind::Worker: {
+    const auto *WC = cast<OpenACCWorkerClause>(C);
+    writeSourceLocation(WC->getLParenLoc());
+    writeBool(WC->hasIntExpr());
+    if (WC->hasIntExpr())
+      AddStmt(const_cast<Expr *>(WC->getIntExpr()));
+    return;
+  }
 
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
-  case OpenACCClauseKind::Worker:
   case OpenACCClauseKind::Vector:
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp
index baa4b173f88edc..ee11435aaa4b1c 100644
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -177,4 +177,43 @@ void foo() {
 #pragma acc serial
 #pragma acc loop gang
   for(;;);
+
+// CHECK: #pragma acc loop worker
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop worker
+  for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop worker
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop worker
+  for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop worker
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop worker
+  for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop worker(num: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop worker(5)
+  for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop worker(num: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop worker(num:5)
+  for(;;);
 }
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 899fbd78b87298..81c48335cf0c42 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -949,7 +949,6 @@ void IntExprParsing() {
 #pragma acc loop vector(length:returns_int())
   for(;;);
 
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc loop worker
   for(;;);
   // expected-error at +1{{expected expression}}
@@ -959,8 +958,8 @@ void IntExprParsing() {
   // expected-error at +1{{expected expression}}
 #pragma acc loop worker(invalid:)
   for(;;);
-  // expected-error at +2{{invalid tag 'invalid' on 'worker' clause}}
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
+#pragma acc kernels
+  // expected-error at +1{{invalid tag 'invalid' on 'worker' clause}}
 #pragma acc loop worker(invalid:5)
   for(;;);
   // expected-error at +1{{expected expression}}
@@ -983,21 +982,21 @@ void IntExprParsing() {
   // expected-note at +1{{to match this '('}}
 #pragma acc loop worker(length:6,4)
   for(;;);
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
+#pragma acc kernels
 #pragma acc loop worker(5)
   for(;;);
-  // expected-error at +2{{invalid tag 'length' on 'worker' clause}}
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
+#pragma acc kernels
+  // expected-error at +1{{invalid tag 'length' on 'worker' clause}}
 #pragma acc loop worker(length:5)
   for(;;);
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
+#pragma acc kernels
 #pragma acc loop worker(num:5)
   for(;;);
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
+#pragma acc kernels
 #pragma acc loop worker(returns_int())
   for(;;);
-  // expected-error at +2{{invalid tag 'length' on 'worker' clause}}
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
+#pragma acc kernels
+  // expected-error at +1{{invalid tag 'length' on 'worker' clause}}
 #pragma acc loop worker(length:returns_int())
   for(;;);
 }
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
index 89000517c43fb5..aaf8b76e1f3dfb 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -61,8 +61,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'auto' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) auto
   while(1);
-  // expected-error at +2{{OpenACC clause 'worker' may not follow a 'device_type' clause in a compute construct}}
-  // expected-note at +1{{previous clause is here}}
+  // expected-error at +1{{OpenACC 'worker' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) worker
   while(1);
   // expected-error at +2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a compute construct}}
diff --git a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
index 6c2c79b02a4131..6a975956f3ff5c 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -43,7 +43,6 @@ void uses() {
   // expected-warning at +1{{OpenACC clause 'if_present' not yet implemented}}
 #pragma acc loop auto if_present
   for(;;);
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
 #pragma acc loop auto worker
   for(;;);
   // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
@@ -180,7 +179,6 @@ void uses() {
   // expected-warning at +1{{OpenACC clause 'if_present' not yet implemented}}
 #pragma acc loop if_present auto
   for(;;);
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
 #pragma acc loop worker auto
   for(;;);
   // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
@@ -318,7 +316,6 @@ void uses() {
   // expected-warning at +1{{OpenACC clause 'if_present' not yet implemented}}
 #pragma acc loop independent if_present
   for(;;);
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
 #pragma acc loop independent worker
   for(;;);
   // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
@@ -455,7 +452,6 @@ void uses() {
   // expected-warning at +1{{OpenACC clause 'if_present' not yet implemented}}
 #pragma acc loop if_present independent
   for(;;);
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
 #pragma acc loop worker independent
   for(;;);
   // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
@@ -591,9 +587,8 @@ void uses() {
   // expected-note at +1{{previous clause is here}}
 #pragma acc loop seq gang
   for(;;);
-  // expected-error at +3{{OpenACC clause 'worker' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
-  // expected-note at +2{{previous clause is here}}
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
+  // expected-error at +2{{OpenACC clause 'worker' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
+  // expected-note at +1{{previous clause is here}}
 #pragma acc loop seq worker
   for(;;);
   // expected-error at +3{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
@@ -734,10 +729,8 @@ void uses() {
   // expected-note at +1{{previous clause is here}}
 #pragma acc loop gang seq
   for(;;);
-  // TODO OpenACC: when 'worker' is implemented and makes it to the AST, this should diagnose because of a conflict with 'seq'.
-  // TODOexpected-error at +3{{OpenACC clause 'worker' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
-  // TODOexpected-note at +2{{previous clause is here}}
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented}}
+  // expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'loop' construct}}
+  // expected-note at +1{{previous clause is here}}
 #pragma acc loop worker seq
   for(;;);
   // TODO OpenACC: when 'vector' is implemented and makes it to the AST, this should diagnose because of a conflict with 'seq'.
diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index cedef3ca858f5e..51da8565f4e399 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -56,7 +56,6 @@ void uses() {
   for(;;);
 #pragma acc loop device_type(*) auto
   for(;;);
-  // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc loop device_type(*) worker
   for(;;);
   // expected-error at +2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'loop' construct}}
diff --git a/clang/test/SemaOpenACC/loop-construct-worker-ast.cpp b/clang/test/SemaOpenACC/loop-construct-worker-ast.cpp
new file mode 100644
index 00000000000000..6347e1419fd5c6
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-worker-ast.cpp
@@ -0,0 +1,270 @@
+// 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
+
+template<unsigned I, typename ConvertsToInt, typename Int>
+void TemplUses(ConvertsToInt CTI, Int IsI) {
+  // CHECK: FunctionTemplateDecl{{.*}}TemplUses
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 0 I
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 ConvertsToInt
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 2 Int
+  // CHECK-NEXT: FunctionDecl{{.*}}TemplUses 'void (ConvertsToInt, Int)'
+  // CHECK-NEXT: ParmVarDecl{{.*}}CTI 'ConvertsToInt'
+  // CHECK-NEXT: ParmVarDecl{{.*}}IsI 'Int'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop worker
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop worker
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop worker
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'ConvertsToInt' lvalue ParmVar
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop worker(CTI)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'Int' lvalue ParmVar
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop worker(num:IsI)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}}'I' 'unsigned int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop worker(num:I)
+  for(;;);
+
+  // Instantiations:
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (Converts, int)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument integral '3U'
+  // CHECK-NEXT: TemplateArgument type 'Converts'
+  // CHECK-NEXT: RecordType{{.*}}'Converts'
+  // CHECK-NEXT: CXXRecord{{.*}}'Converts
+  // CHECK-NEXT: TemplateArgument type 'int'
+  // CHECK-NEXT: BuiltinType{{.*}}'int'
+  // CHECK-NEXT: ParmVarDecl{{.*}} CTI 'Converts'
+  // CHECK-NEXT: ParmVarDecl{{.*}} IsI 'int'
+  // CHECK-NEXT: CompoundStmt
+  //
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'Converts' lvalue ParmVar
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'IsI' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}}'unsigned int' depth 0 index 0 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+}
+
+struct Converts{
+  operator int();
+};
+
+void uses() {
+  // CHECK: FunctionDecl{{.*}} uses
+  // CHECK-NEXT: CompoundStmt
+  //
+  // CHECK-NEXT: CallExpr
+  TemplUses<3>(Converts{}, 5);
+
+  // CHECK: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop worker
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop worker
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop worker
+  for(;;);
+
+  Converts CTI;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+  // CHECK-NEXT: CXXConstructExpr
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue Var
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop worker(CTI)
+  for(;;);
+
+  int IsI;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: worker clause{{.*}}
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop worker(num:IsI)
+  for(;;);
+}
+
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/loop-construct-worker-clause.cpp b/clang/test/SemaOpenACC/loop-construct-worker-clause.cpp
new file mode 100644
index 00000000000000..f7d2f365a5aa79
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-worker-clause.cpp
@@ -0,0 +1,202 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+template<unsigned I, typename NotInt, typename ConvertsToInt, typename Int>
+void TemplUses(NotInt NI, ConvertsToInt CTI, Int IsI) {
+  int i;
+
+  // expected-error at +1{{'num' argument on 'worker' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop worker(i)
+  for(;;);
+
+  // expected-error at +1{{'num' argument on 'worker' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop worker(num:IsI)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(i)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(CTI)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(IsI)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(I)
+  for(;;);
+
+#pragma acc kernels
+  // expected-error at +1{{OpenACC clause 'worker' requires expression of integer type ('NoConvert' invalid)}}
+#pragma acc loop worker(NI)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(num:i)
+  for(;;);
+
+  // expected-error at +3{{'num' argument to 'worker' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_workers' clause}}
+  // expected-note at +1{{previous clause is here}}
+#pragma acc kernels num_workers(IsI)
+#pragma acc loop worker(num:CTI)
+  for(;;);
+ for(;;);
+}
+
+struct NoConvert{};
+struct Converts{
+  operator int();
+};
+
+void uses() {
+  TemplUses<3>(NoConvert{}, Converts{}, 5); // expected-note{{in instantiation of function template specialization}}
+
+#pragma acc loop worker
+  for(;;);
+
+#pragma acc parallel
+#pragma acc loop worker
+  for(;;);
+
+  int i;
+
+  // expected-error at +1{{'num' argument on 'worker' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop worker(i)
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'worker' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop worker(i)
+  for(;;);
+
+  // expected-error at +1{{'num' argument on 'worker' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop worker(num:i)
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'worker' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop worker(num:i)
+  for(;;);
+
+#pragma acc serial
+#pragma acc loop worker
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'worker' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop worker(i)
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'worker' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop worker(num:i)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(i)
+  for(;;);
+
+  Converts Cvts;
+
+#pragma acc kernels
+#pragma acc loop worker(Cvts)
+  for(;;);
+
+  NoConvert NoCvts;
+
+#pragma acc kernels
+  // expected-error at +1{{OpenACC clause 'worker' requires expression of integer type ('NoConvert' invalid)}}
+#pragma acc loop worker(NoCvts)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(num:i)
+  for(;;);
+
+  // expected-error at +3{{'num' argument to 'worker' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_workers' clause}}
+  // expected-note at +1{{previous clause is here}}
+#pragma acc kernels num_workers(i)
+#pragma acc loop worker(num:i)
+  for(;;);
+
+#pragma acc loop worker
+  for(;;) {
+    // expected-error at +3{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}}
+    // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}}
+    // expected-note at -4 2{{previous clause is here}}
+#pragma acc loop worker, gang
+  for(;;) {}
+  }
+
+#pragma acc loop worker
+  for(;;) {
+#pragma acc parallel
+#pragma acc loop worker, gang
+  for(;;) {}
+  }
+
+
+#pragma acc parallel
+#pragma acc loop worker
+  for(;;) {
+    // expected-error at +3{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}}
+    // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}}
+    // expected-note at -4 2{{previous clause is here}}
+#pragma acc loop worker, gang
+  for(;;) {}
+  }
+
+#pragma acc parallel
+#pragma acc loop worker
+  for(;;) {
+#pragma acc parallel
+#pragma acc loop worker, gang
+  for(;;) {}
+  }
+
+#pragma acc serial
+#pragma acc loop worker
+  for(;;) {
+    // expected-error at +3{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}}
+    // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}}
+    // expected-note at -4 2{{previous clause is here}}
+#pragma acc loop worker, gang
+  for(;;) {}
+  }
+
+#pragma acc serial
+#pragma acc loop worker
+  for(;;) {
+#pragma acc parallel
+#pragma acc loop worker, gang
+  for(;;) {}
+  }
+
+#pragma acc kernels
+#pragma acc loop worker
+  for(;;) {
+    // expected-error at +3{{loop with a 'worker' clause may not exist in the region of a 'worker' clause}}
+    // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'worker' clause}}
+    // expected-note at -4 2{{previous clause is here}}
+#pragma acc loop worker, gang
+  for(;;) {}
+  }
+
+#pragma acc kernels
+#pragma acc loop worker
+  for(;;) {
+#pragma acc parallel
+#pragma acc loop worker, gang
+  for(;;) {}
+  }
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 2ffe47fbd74476..4461be86ea9996 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2888,6 +2888,11 @@ void OpenACCClauseEnqueue::VisitAsyncClause(const OpenACCAsyncClause &C) {
   if (C.hasIntExpr())
     Visitor.AddStmt(C.getIntExpr());
 }
+
+void OpenACCClauseEnqueue::VisitWorkerClause(const OpenACCWorkerClause &C) {
+  if (C.hasIntExpr())
+    Visitor.AddStmt(C.getIntExpr());
+}
 void OpenACCClauseEnqueue::VisitWaitClause(const OpenACCWaitClause &C) {
   if (const Expr *DevNumExpr = C.getDevNumExpr())
     Visitor.AddStmt(DevNumExpr);



More information about the cfe-commits mailing list