[clang] [OpenACC] Implement 'num_workers' clause for compute constructs (PR #89151)

Erich Keane via cfe-commits cfe-commits at lists.llvm.org
Wed Apr 17 15:18:08 PDT 2024


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

This clause just takes an 'int expr', which is not optional.  This patch implements the clause on compute constructs.

>From d3894971090921b92c71ba5a18151cb2033c8cfa Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Tue, 16 Apr 2024 09:43:55 -0700
Subject: [PATCH] [OpenACC] Implement 'num_workers' clause for compute
 constructs

This clause just takes an 'int expr', which is not optional.  This patch
implements the clause on compute constructs.
---
 clang/include/clang/AST/OpenACCClause.h       |  56 ++++
 .../clang/Basic/DiagnosticSemaKinds.td        |  12 +
 clang/include/clang/Basic/OpenACCClauses.def  |   1 +
 clang/include/clang/Parse/Parser.h            |   9 +-
 clang/include/clang/Sema/SemaOpenACC.h        |  36 ++-
 clang/lib/AST/OpenACCClause.cpp               |  26 ++
 clang/lib/AST/StmtProfile.cpp                 |   7 +
 clang/lib/AST/TextNodeDumper.cpp              |   1 +
 clang/lib/Parse/ParseOpenACC.cpp              |  57 ++--
 clang/lib/Sema/SemaOpenACC.cpp                | 122 +++++++++
 clang/lib/Sema/TreeTransform.h                |  23 ++
 clang/lib/Serialization/ASTReader.cpp         |   7 +-
 clang/lib/Serialization/ASTWriter.cpp         |   7 +-
 clang/test/ParserOpenACC/parse-clauses.c      |   2 -
 .../compute-construct-intexpr-clause-ast.cpp  | 255 ++++++++++++++++++
 .../compute-construct-num_workers-clause.c    |  33 +++
 .../compute-construct-num_workers-clause.cpp  | 133 +++++++++
 clang/tools/libclang/CIndex.cpp               |   4 +
 18 files changed, 764 insertions(+), 27 deletions(-)
 create mode 100644 clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
 create mode 100644 clang/test/SemaOpenACC/compute-construct-num_workers-clause.c
 create mode 100644 clang/test/SemaOpenACC/compute-construct-num_workers-clause.cpp

diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 07587849eb1219..7a60620d5875c5 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -156,6 +156,62 @@ class OpenACCSelfClause : public OpenACCClauseWithCondition {
                                    Expr *ConditionExpr, SourceLocation EndLoc);
 };
 
+/// Represents one of a handful of classes that have integer expressions.
+/// Semantically, many only permit a single expression, with a few that permit
+/// up to 3.
+class OpenACCClauseWithIntExprs : public OpenACCClauseWithParams {
+  llvm::SmallVector<Expr *> IntExprs;
+
+  protected:
+    OpenACCClauseWithIntExprs(OpenACCClauseKind K, SourceLocation BeginLoc,
+                              SourceLocation LParenLoc,
+                              ArrayRef<Expr *> IntExprs, SourceLocation EndLoc)
+        : OpenACCClauseWithParams(K, BeginLoc, LParenLoc, EndLoc),
+          IntExprs(IntExprs) {}
+
+    /// Gets the entire list of integer expressions, but leave it to the
+    /// individual clauses to expose this how they'd like.
+    llvm::ArrayRef<Expr *> getIntExprs() const { return IntExprs; }
+
+  public:
+  child_range children() {
+    return child_range(reinterpret_cast<Stmt **>(IntExprs.begin()),
+                       reinterpret_cast<Stmt **>(IntExprs.end()));
+  }
+
+  const_child_range children() const {
+    child_range Children =
+        const_cast<OpenACCClauseWithIntExprs *>(this)->children();
+    return const_child_range(Children.begin(), Children.end());
+  }
+};
+
+/// A more restrictive version of the IntExprs version that exposes a single
+/// integer expression.
+class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithIntExprs {
+  protected:
+    OpenACCClauseWithSingleIntExpr(OpenACCClauseKind K, SourceLocation BeginLoc,
+                                   SourceLocation LParenLoc, Expr *IntExpr,
+                                   SourceLocation EndLoc)
+        : OpenACCClauseWithIntExprs(K, BeginLoc, LParenLoc, IntExpr, EndLoc) {}
+
+  public:
+    bool hasIntExpr() const { return !getIntExprs().empty(); }
+    const Expr *getIntExpr() const {
+      return hasIntExpr() ? getIntExprs()[0] : nullptr;
+    }
+    Expr *getIntExpr() { return hasIntExpr() ? getIntExprs()[0] : nullptr; }
+};
+
+class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr {
+  OpenACCNumWorkersClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                          Expr *IntExpr, SourceLocation EndLoc);
+  public:
+    static OpenACCNumWorkersClause *
+    Create(const ASTContext &C, SourceLocation BeginLoc,
+           SourceLocation LParenLoc, Expr *IntExpr, SourceLocation EndLoc);
+};
+
 template <class Impl> class OpenACCClauseVisitor {
   Impl &getDerived() { return static_cast<Impl &>(*this); }
 
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 30a8543489f48e..5ac1b3dc6233a3 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12268,4 +12268,16 @@ def warn_acc_if_self_conflict
     : Warning<"OpenACC construct 'self' has no effect when an 'if' clause "
               "evaluates to true">,
       InGroup<DiagGroup<"openacc-self-if-potential-conflict">>;
+def err_acc_int_expr_requires_integer
+    : Error<"OpenACC %select{clause|directive}0 '%1' requires expression of "
+            "integer type (%2 invalid)">;
+def err_acc_int_expr_incomplete_class_type
+    : Error<"OpenACC integer expression has incomplete class type %0">;
+def err_acc_int_expr_explicit_conversion
+    : Error<"OpenACC integer expression type %0 requires explicit conversion "
+            "to %1">;
+def note_acc_int_expr_conversion
+    : Note<"conversion to %select{integral|enumeration}0 type %1">;
+def err_acc_int_expr_multiple_conversions
+    : Error<"multiple conversions from expression type %0 to an integral type">;
 } // end of sema component.
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 378495d2c0909a..d1a95cbe613944 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -18,5 +18,6 @@
 VISIT_CLAUSE(Default)
 VISIT_CLAUSE(If)
 VISIT_CLAUSE(Self)
+VISIT_CLAUSE(NumWorkers)
 
 #undef VISIT_CLAUSE
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 23b268126de4e0..72b2f958a5e622 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3640,13 +3640,14 @@ class Parser : public CodeCompletionHandler {
   /// Parses the clause-list for an OpenACC directive.
   SmallVector<OpenACCClause *>
   ParseOpenACCClauseList(OpenACCDirectiveKind DirKind);
-  bool ParseOpenACCWaitArgument();
+  bool ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective);
   /// Parses the clause of the 'bind' argument, which can be a string literal or
   /// an ID expression.
   ExprResult ParseOpenACCBindClauseArgument();
   /// Parses the clause kind of 'int-expr', which can be any integral
   /// expression.
-  ExprResult ParseOpenACCIntExpr();
+  ExprResult ParseOpenACCIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+                                 SourceLocation Loc);
   /// Parses the 'device-type-list', which is a list of identifiers.
   bool ParseOpenACCDeviceTypeList();
   /// Parses the 'async-argument', which is an integral value with two
@@ -3657,9 +3658,9 @@ class Parser : public CodeCompletionHandler {
   /// Parses a comma delimited list of 'size-expr's.
   bool ParseOpenACCSizeExprList();
   /// Parses a 'gang-arg-list', used for the 'gang' clause.
-  bool ParseOpenACCGangArgList();
+  bool ParseOpenACCGangArgList(SourceLocation GangLoc);
   /// Parses a 'gang-arg', used for the 'gang' clause.
-  bool ParseOpenACCGangArg();
+  bool ParseOpenACCGangArg(SourceLocation GangLoc);
   /// Parses a 'condition' expr, ensuring it results in a
   ExprResult ParseOpenACCConditionExpr();
 
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 329dc3945fa2a6..eb461fa7dbd541 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -44,8 +44,13 @@ class SemaOpenACC : public SemaBase {
       Expr *ConditionExpr;
     };
 
-    std::variant<std::monostate, DefaultDetails, ConditionDetails> Details =
-        std::monostate{};
+    struct IntExprDetails {
+      SmallVector<Expr *> IntExprs;
+    };
+
+    std::variant<std::monostate, DefaultDetails, ConditionDetails,
+                 IntExprDetails>
+        Details = std::monostate{};
 
   public:
     OpenACCParsedClause(OpenACCDirectiveKind DirKind,
@@ -87,6 +92,22 @@ class SemaOpenACC : public SemaBase {
       return std::get<ConditionDetails>(Details).ConditionExpr;
     }
 
+    unsigned getNumIntExprs() const {
+      assert(ClauseKind == OpenACCClauseKind::NumWorkers &&
+             "Parsed clause kind does not have a int exprs");
+      return std::get<IntExprDetails>(Details).IntExprs.size();
+    }
+
+    ArrayRef<Expr *> getIntExprs() {
+      assert(ClauseKind == OpenACCClauseKind::NumWorkers &&
+             "Parsed clause kind does not have a int exprs");
+      return std::get<IntExprDetails>(Details).IntExprs;
+    }
+
+    ArrayRef<Expr *> getIntExprs() const {
+      return const_cast<OpenACCParsedClause*>(this)->getIntExprs();
+    }
+
     void setLParenLoc(SourceLocation EndLoc) { LParenLoc = EndLoc; }
     void setEndLoc(SourceLocation EndLoc) { ClauseRange.setEnd(EndLoc); }
 
@@ -109,6 +130,12 @@ class SemaOpenACC : public SemaBase {
 
       Details = ConditionDetails{ConditionExpr};
     }
+
+    void setIntExprDetails(ArrayRef<Expr *> IntExprs) {
+      assert(ClauseKind == OpenACCClauseKind::NumWorkers &&
+             "Parsed clause kind does not have a int exprs");
+      Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}};
+    }
   };
 
   SemaOpenACC(Sema &S);
@@ -148,6 +175,11 @@ class SemaOpenACC : public SemaBase {
   /// Called after the directive has been completely parsed, including the
   /// declaration group or associated statement.
   DeclGroupRef ActOnEndDeclDirective();
+
+  /// Called when encountering an 'int-expr' for OpenACC, and manages
+  /// conversions and diagnostics to 'int'.
+  ExprResult ActOnIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+                          SourceLocation Loc, Expr *IntExpr);
 };
 
 } // namespace clang
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 9c259c8f9bd0a1..3fff02fa33f28d 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -82,6 +82,27 @@ OpenACCClause::child_range OpenACCClause::children() {
   return child_range(child_iterator(), child_iterator());
 }
 
+OpenACCNumWorkersClause::OpenACCNumWorkersClause(SourceLocation BeginLoc,
+                                                 SourceLocation LParenLoc,
+                                                 Expr *IntExpr,
+                                                 SourceLocation EndLoc)
+    : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::NumWorkers, BeginLoc,
+                                     LParenLoc, IntExpr, EndLoc) {
+  assert((!IntExpr || IntExpr->isInstantiationDependent() ||
+          IntExpr->getType()->isIntegerType()) &&
+         "Condition expression type not scalar/dependent");
+}
+
+OpenACCNumWorkersClause *
+OpenACCNumWorkersClause::Create(const ASTContext &C, SourceLocation BeginLoc,
+                                SourceLocation LParenLoc, Expr *IntExpr,
+                                SourceLocation EndLoc) {
+  void *Mem = C.Allocate(sizeof(OpenACCNumWorkersClause),
+                         alignof(OpenACCNumWorkersClause));
+  return new (Mem)
+      OpenACCNumWorkersClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
+}
+
 //===----------------------------------------------------------------------===//
 //  OpenACC clauses printing methods
 //===----------------------------------------------------------------------===//
@@ -98,3 +119,8 @@ void OpenACCClausePrinter::VisitSelfClause(const OpenACCSelfClause &C) {
   if (const Expr *CondExpr = C.getConditionExpr())
     OS << "(" << CondExpr << ")";
 }
+
+void OpenACCClausePrinter::VisitNumWorkersClause(
+    const OpenACCNumWorkersClause &C) {
+  OS << "num_workers(" << C.getIntExpr() << ")";
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index b26d804c6f079b..ab7d4c5f930112 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2496,6 +2496,13 @@ void OpenACCClauseProfiler::VisitSelfClause(const OpenACCSelfClause &Clause) {
   if (Clause.hasConditionExpr())
     Profiler.VisitStmt(Clause.getConditionExpr());
 }
+
+void OpenACCClauseProfiler::VisitNumWorkersClause(
+    const OpenACCNumWorkersClause &Clause) {
+  assert(Clause.hasIntExpr() && "if clause requires a valid condition expr");
+  Profiler.VisitStmt(Clause.getIntExpr());
+}
+
 } // namespace
 
 void StmtProfiler::VisitOpenACCComputeConstruct(
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index ff5b3df2d6dfac..9d1b73cb7a0784 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -399,6 +399,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
       break;
     case OpenACCClauseKind::If:
     case OpenACCClauseKind::Self:
+    case OpenACCClauseKind::NumWorkers:
       // 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.
       OS << " clause";
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 123be476e928ee..096e0863ed47c7 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -632,10 +632,16 @@ Parser::ParseOpenACCClauseList(OpenACCDirectiveKind DirKind) {
   return Clauses;
 }
 
-ExprResult Parser::ParseOpenACCIntExpr() {
-  // FIXME: this is required to be an integer expression (or dependent), so we
-  // should ensure that is the case by passing this to SEMA here.
-  return getActions().CorrectDelayedTyposInExpr(ParseAssignmentExpression());
+ExprResult Parser::ParseOpenACCIntExpr(OpenACCDirectiveKind DK,
+                                       OpenACCClauseKind CK,
+                                       SourceLocation Loc) {
+  ExprResult ER =
+      getActions().CorrectDelayedTyposInExpr(ParseAssignmentExpression());
+
+  if (!ER.isUsable())
+    return ER;
+
+  return getActions().OpenACC().ActOnIntExpr(DK, CK, Loc, ER.get());
 }
 
 bool Parser::ParseOpenACCClauseVarList(OpenACCClauseKind Kind) {
@@ -739,7 +745,7 @@ bool Parser::ParseOpenACCSizeExprList() {
 /// [num:]int-expr
 /// dim:int-expr
 /// static:size-expr
-bool Parser::ParseOpenACCGangArg() {
+bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
 
   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Static, getCurToken()) &&
       NextToken().is(tok::colon)) {
@@ -753,7 +759,9 @@ bool Parser::ParseOpenACCGangArg() {
       NextToken().is(tok::colon)) {
     ConsumeToken();
     ConsumeToken();
-    return ParseOpenACCIntExpr().isInvalid();
+    return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
+                               OpenACCClauseKind::Gang, GangLoc)
+        .isInvalid();
   }
 
   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Num, getCurToken()) &&
@@ -763,11 +771,13 @@ bool Parser::ParseOpenACCGangArg() {
     // Fallthrough to the 'int-expr' handling for when 'num' is omitted.
   }
   // This is just the 'num' case where 'num' is optional.
-  return ParseOpenACCIntExpr().isInvalid();
+  return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
+                             OpenACCClauseKind::Gang, GangLoc)
+      .isInvalid();
 }
 
-bool Parser::ParseOpenACCGangArgList() {
-  if (ParseOpenACCGangArg()) {
+bool Parser::ParseOpenACCGangArgList(SourceLocation GangLoc) {
+  if (ParseOpenACCGangArg(GangLoc)) {
     SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
               Parser::StopBeforeMatch);
     return false;
@@ -776,7 +786,7 @@ bool Parser::ParseOpenACCGangArgList() {
   while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
     ExpectAndConsume(tok::comma);
 
-    if (ParseOpenACCGangArg()) {
+    if (ParseOpenACCGangArg(GangLoc)) {
       SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
                 Parser::StopBeforeMatch);
       return false;
@@ -941,11 +951,18 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
     case OpenACCClauseKind::DeviceNum:
     case OpenACCClauseKind::DefaultAsync:
     case OpenACCClauseKind::VectorLength: {
-      ExprResult IntExpr = ParseOpenACCIntExpr();
+      ExprResult IntExpr = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
+                                               ClauseKind, ClauseLoc);
       if (IntExpr.isInvalid()) {
         Parens.skipToEnd();
         return OpenACCCanContinue();
       }
+
+      // TODO OpenACC: as we implement the 'rest' of the above, this 'if' should
+      // be removed leaving just the 'setIntExprDetails'.
+      if (ClauseKind == OpenACCClauseKind::NumWorkers)
+        ParsedClause.setIntExprDetails(IntExpr.get());
+
       break;
     }
     case OpenACCClauseKind::DType:
@@ -998,7 +1015,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
                                                ? OpenACCSpecialTokenKind::Length
                                                : OpenACCSpecialTokenKind::Num,
                                            ClauseKind);
-        ExprResult IntExpr = ParseOpenACCIntExpr();
+        ExprResult IntExpr = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
+                                                 ClauseKind, ClauseLoc);
         if (IntExpr.isInvalid()) {
           Parens.skipToEnd();
           return OpenACCCanContinue();
@@ -1014,13 +1032,14 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
         break;
       }
       case OpenACCClauseKind::Gang:
-        if (ParseOpenACCGangArgList()) {
+        if (ParseOpenACCGangArgList(ClauseLoc)) {
           Parens.skipToEnd();
           return OpenACCCanContinue();
         }
         break;
       case OpenACCClauseKind::Wait:
-        if (ParseOpenACCWaitArgument()) {
+        if (ParseOpenACCWaitArgument(ClauseLoc,
+                                     /*IsDirective=*/false)) {
           Parens.skipToEnd();
           return OpenACCCanContinue();
         }
@@ -1052,7 +1071,7 @@ ExprResult Parser::ParseOpenACCAsyncArgument() {
 /// In this section and throughout the specification, the term wait-argument
 /// means:
 /// [ devnum : int-expr : ] [ queues : ] async-argument-list
-bool Parser::ParseOpenACCWaitArgument() {
+bool Parser::ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective) {
   // [devnum : int-expr : ]
   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::DevNum, Tok) &&
       NextToken().is(tok::colon)) {
@@ -1061,7 +1080,11 @@ bool Parser::ParseOpenACCWaitArgument() {
     // Consume colon.
     ConsumeToken();
 
-    ExprResult IntExpr = ParseOpenACCIntExpr();
+    ExprResult IntExpr = ParseOpenACCIntExpr(
+        IsDirective ? OpenACCDirectiveKind::Wait
+                    : OpenACCDirectiveKind::Invalid,
+        IsDirective ? OpenACCClauseKind::Invalid : OpenACCClauseKind::Wait,
+        Loc);
     if (IntExpr.isInvalid())
       return true;
 
@@ -1245,7 +1268,7 @@ Parser::OpenACCDirectiveParseInfo Parser::ParseOpenACCDirective() {
       break;
     case OpenACCDirectiveKind::Wait:
       // OpenACC has an optional paren-wrapped 'wait-argument'.
-      if (ParseOpenACCWaitArgument())
+      if (ParseOpenACCWaitArgument(StartLoc, /*IsDirective=*/true))
         T.skipToEnd();
       else
         T.consumeClose();
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 59f65eaf47a6da..316b0f15b049be 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -14,6 +14,7 @@
 #include "clang/Sema/SemaOpenACC.h"
 #include "clang/AST/StmtOpenACC.h"
 #include "clang/Basic/DiagnosticSema.h"
+#include "clang/Basic/OpenACCKinds.h"
 #include "clang/Sema/Sema.h"
 #include "llvm/Support/Casting.h"
 
@@ -90,6 +91,17 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
     default:
       return false;
     }
+  case OpenACCClauseKind::NumWorkers:
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Parallel:
+    case OpenACCDirectiveKind::Kernels:
+    case OpenACCDirectiveKind::Update:
+    case OpenACCDirectiveKind::ParallelLoop:
+    case OpenACCDirectiveKind::KernelsLoop:
+      return true;
+    default:
+      return false;
+    }
   default:
     // Do nothing so we can go to the 'unimplemented' diagnostic instead.
     return true;
@@ -218,6 +230,25 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
         getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
         Clause.getConditionExpr(), Clause.getEndLoc());
   }
+  case OpenACCClauseKind::NumWorkers: {
+    // Restrictions only properly implemented on 'compute' constructs, and
+    // 'compute' constructs are the only construct that can do anything with
+    // this yet, so skip/treat as unimplemented in this case.
+    if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+      break;
+
+    // There is no prose in the standard that says duplicates aren't allowed,
+    // but this diagnostic is present in other compilers, as well as makes
+    // sense.
+    if (checkAlreadyHasClauseOfKind(*this, ExistingClauses, Clause))
+      return nullptr;
+
+    assert(Clause.getIntExprs().size() == 1 &&
+           "Invalid number of expressions for NumWorkers");
+    return OpenACCNumWorkersClause::Create(
+        getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
+        Clause.getIntExprs()[0], Clause.getEndLoc());
+  }
   default:
     break;
   }
@@ -248,6 +279,97 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K,
   }
 }
 
+ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK,
+                                     OpenACCClauseKind CK,
+                                     SourceLocation Loc, Expr *IntExpr) {
+
+  assert(((DK != OpenACCDirectiveKind::Invalid &&
+           CK == OpenACCClauseKind::Invalid) ||
+          (DK == OpenACCDirectiveKind::Invalid &&
+           CK != OpenACCClauseKind::Invalid)) &&
+         "Only one of directive or clause kind should be provided");
+
+  class IntExprConverter : public Sema::ICEConvertDiagnoser {
+    OpenACCDirectiveKind DirectiveKind;
+    OpenACCClauseKind ClauseKind;
+    Expr *IntExpr;
+
+  public:
+    IntExprConverter(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+                     Expr *IntExpr)
+        : ICEConvertDiagnoser(/*AllowScopedEnumerations=*/false,
+                              /*Suppress=*/false,
+                              /*SuppressConversion=*/true),
+          DirectiveKind(DK), ClauseKind(CK), IntExpr(IntExpr) {}
+
+    bool match(QualType T) override {
+      // OpenACC spec just calls this 'integer expression' as having an
+      // 'integer type', so fall back on C99's 'integer type'.
+      return T->isIntegerType();
+    }
+    SemaBase::SemaDiagnosticBuilder diagnoseNotInt(Sema &S, SourceLocation Loc,
+                                                   QualType T) override {
+      if (ClauseKind != OpenACCClauseKind::Invalid)
+        return S.Diag(Loc, diag::err_acc_int_expr_requires_integer) <<
+               /*Clause=*/0 << ClauseKind << T;
+
+      return S.Diag(Loc, diag::err_acc_int_expr_requires_integer) <<
+             /*Directive=*/1 << DirectiveKind << T;
+    }
+
+    SemaBase::SemaDiagnosticBuilder
+    diagnoseIncomplete(Sema &S, SourceLocation Loc, QualType T) override {
+      return S.Diag(Loc, diag::err_acc_int_expr_incomplete_class_type)
+             << T << IntExpr->getSourceRange();
+    }
+
+    SemaBase::SemaDiagnosticBuilder
+    diagnoseExplicitConv(Sema &S, SourceLocation Loc, QualType T,
+                         QualType ConvTy) override {
+      return S.Diag(Loc, diag::err_acc_int_expr_explicit_conversion)
+             << T << ConvTy;
+    }
+
+    SemaBase::SemaDiagnosticBuilder noteExplicitConv(Sema &S,
+                                                     CXXConversionDecl *Conv,
+                                                     QualType ConvTy) override {
+      return S.Diag(Conv->getLocation(), diag::note_acc_int_expr_conversion)
+             << ConvTy->isEnumeralType() << ConvTy;
+    }
+
+    SemaBase::SemaDiagnosticBuilder
+    diagnoseAmbiguous(Sema &S, SourceLocation Loc, QualType T) override {
+      return S.Diag(Loc, diag::err_acc_int_expr_multiple_conversions) << T;
+    }
+
+    SemaBase::SemaDiagnosticBuilder
+    noteAmbiguous(Sema &S, CXXConversionDecl *Conv, QualType ConvTy) override {
+      return S.Diag(Conv->getLocation(), diag::note_acc_int_expr_conversion)
+             << ConvTy->isEnumeralType() << ConvTy;
+    }
+
+    SemaBase::SemaDiagnosticBuilder
+    diagnoseConversion(Sema &S, SourceLocation Loc, QualType T,
+                       QualType ConvTy) override {
+      llvm_unreachable("conversion functions are permitted");
+    }
+  } IntExprDiagnoser(DK, CK, IntExpr);
+
+  ExprResult IntExprResult = SemaRef.PerformContextualImplicitConversion(
+      Loc, IntExpr, IntExprDiagnoser);
+  if (IntExprResult.isInvalid())
+    return ExprError();
+
+  IntExpr = IntExprResult.get();
+  if (!IntExpr->isTypeDependent() &&
+      !IntExpr->getType()->isIntegerType())
+    return ExprError();
+
+  // TODO OpenACC: Do we want to perform usual unary conversions here? When
+  // doing codegen we might find that is necessary, but skip it for now.
+  return IntExpr;
+}
+
 bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
                                           SourceLocation StartLoc) {
   return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/true);
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index eb05783a6219dc..34b52a346a10d1 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11158,6 +11158,29 @@ void OpenACCClauseTransform<Derived>::VisitSelfClause(
       ParsedClause.getLParenLoc(), ParsedClause.getConditionExpr(),
       ParsedClause.getEndLoc());
 }
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitNumWorkersClause(
+    const OpenACCNumWorkersClause &C) {
+  Expr *IntExpr = const_cast<Expr *>(C.getIntExpr());
+  assert(IntExpr && "num_workers clause constructed with invalid int expr");
+
+  ExprResult Res = Self.TransformExpr(IntExpr);
+  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 = OpenACCNumWorkersClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getIntExprs()[0],
+      ParsedClause.getEndLoc());
+}
 } // namespace
 template <typename Derived>
 OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index b28df03b4a95e9..6fff829cefb845 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11786,6 +11786,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCSelfClause::Create(getContext(), BeginLoc, LParenLoc,
                                      CondExpr, EndLoc);
   }
+  case OpenACCClauseKind::NumWorkers: {
+    SourceLocation LParenLoc = readSourceLocation();
+    Expr *IntExpr = readSubExpr();
+    return OpenACCNumWorkersClause::Create(getContext(), BeginLoc, LParenLoc,
+                                           IntExpr, EndLoc);
+  }
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::Seq:
@@ -11816,7 +11822,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   case OpenACCClauseKind::Bind:
   case OpenACCClauseKind::VectorLength:
   case OpenACCClauseKind::NumGangs:
-  case OpenACCClauseKind::NumWorkers:
   case OpenACCClauseKind::DeviceNum:
   case OpenACCClauseKind::DefaultAsync:
   case OpenACCClauseKind::DeviceType:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index b2a078b6d80f46..62b81342b052e4 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -7532,6 +7532,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
       AddStmt(const_cast<Expr *>(SC->getConditionExpr()));
     return;
   }
+  case OpenACCClauseKind::NumWorkers: {
+    const auto *NWC = cast<OpenACCNumWorkersClause>(C);
+    writeSourceLocation(NWC->getLParenLoc());
+    AddStmt(const_cast<Expr *>(NWC->getIntExpr()));
+    return;
+  }
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::Seq:
@@ -7562,7 +7568,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::Bind:
   case OpenACCClauseKind::VectorLength:
   case OpenACCClauseKind::NumGangs:
-  case OpenACCClauseKind::NumWorkers:
   case OpenACCClauseKind::DeviceNum:
   case OpenACCClauseKind::DefaultAsync:
   case OpenACCClauseKind::DeviceType:
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 4462f0df540f2d..bd5f459eae074f 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -943,11 +943,9 @@ void IntExprParsing() {
 #pragma acc parallel num_workers(5, 4)
   {}
 
-  // expected-warning at +1{{OpenACC clause 'num_workers' not yet implemented, clause ignored}}
 #pragma acc parallel num_workers(5)
   {}
 
-  // expected-warning at +1{{OpenACC clause 'num_workers' not yet implemented, clause ignored}}
 #pragma acc parallel num_workers(returns_int())
   {}
 
diff --git a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
new file mode 100644
index 00000000000000..e3664460a6e5e3
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
@@ -0,0 +1,255 @@
+// 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
+
+int some_int();
+short some_short();
+long some_long();
+enum E{};
+E some_enum();
+
+struct CorrectConvert {
+  operator int();
+} Convert;
+
+
+void NormalUses() {
+  // CHECK: FunctionDecl{{.*}}NormalUses
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_workers(some_int())
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels num_workers(some_short())
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CallExpr{{.*}}'short'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_workers(some_long())
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CallExpr{{.*}}'long'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'long (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'long ()' lvalue Function{{.*}} 'some_long' 'long ()'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_workers(some_enum())
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CallExpr{{.*}}'E'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'E (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'E ()' lvalue Function{{.*}} 'some_enum' 'E ()'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels num_workers(Convert)
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'struct CorrectConvert':'CorrectConvert' lvalue Var
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+}
+
+template<typename T, typename U>
+void TemplUses(T t, U u) {
+  // CHECK-NEXT: FunctionTemplateDecl
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 U
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T, U)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} referenced t 'T'
+  // CHECK-NEXT: ParmVarDecl{{.*}} referenced u 'U'
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_workers(t)
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T' lvalue ParmVar{{.*}} 't' 'T'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels num_workers(u)
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_workers(U::value)
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels num_workers(T{})
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'T' 'T' list
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_workers(U{})
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'U' 'U' list
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels num_workers(typename U::IntTy{})
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename U::IntTy' 'typename U::IntTy' list
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_workers(typename U::ShortTy{})
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename U::ShortTy' 'typename U::ShortTy' list
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+  // Check the instantiated versions of the above.
+  // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'CorrectConvert'
+  // CHECK-NEXT: RecordType{{.*}} 'CorrectConvert'
+  // CHECK-NEXT: CXXRecord{{.*}} 'CorrectConvert'
+  // CHECK-NEXT: TemplateArgument type 'HasInt'
+  // CHECK-NEXT: RecordType{{.*}} 'HasInt'
+  // CHECK-NEXT: CXXRecord{{.*}} 'HasInt'
+  // CHECK-NEXT: ParmVarDecl{{.*}} used t 'CorrectConvert'
+  // CHECK-NEXT: ParmVarDecl{{.*}} used u 'HasInt'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'CorrectConvert' lvalue ParmVar
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator int
+  // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'CorrectConvert' lvalue
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'CorrectConvert' functional cast to struct CorrectConvert <NoOp>
+  // CHECK-NEXT: InitListExpr{{.*}}'CorrectConvert'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: ExprWithCleanups
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
+  // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'HasInt' lvalue
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'HasInt' functional cast to struct HasInt <NoOp>
+  // CHECK-NEXT: InitListExpr{{.*}}'HasInt'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: ExprWithCleanups
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'typename HasInt::IntTy':'int' functional cast to typename struct HasInt::IntTy <NoOp>
+  // CHECK-NEXT: InitListExpr{{.*}}'typename HasInt::IntTy':'int'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'typename HasInt::ShortTy':'short' functional cast to typename struct HasInt::ShortTy <NoOp>
+  // CHECK-NEXT: InitListExpr{{.*}}'typename HasInt::ShortTy':'short'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+}
+
+struct HasInt {
+  using IntTy = int;
+  using ShortTy = short;
+  static constexpr int value = 1;
+
+  operator char();
+};
+
+void Inst() {
+  TemplUses<CorrectConvert, HasInt>({}, {});
+}
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c b/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c
new file mode 100644
index 00000000000000..19e247a2f810ae
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c
@@ -0,0 +1,33 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+short getS();
+
+void Test() {
+#pragma acc parallel num_workers(1)
+  while(1);
+#pragma acc kernels num_workers(1)
+  while(1);
+
+  // expected-error at +1{{OpenACC 'num_workers' clause is not valid on 'serial' directive}}
+#pragma acc serial num_workers(1)
+  while(1);
+
+  struct NotConvertible{} NC;
+  // expected-error at +1{{OpenACC clause 'num_workers' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_workers(NC)
+  while(1);
+
+#pragma acc kernels num_workers(getS())
+  while(1);
+
+  struct Incomplete *SomeIncomplete;
+
+  // expected-error at +1{{OpenACC clause 'num_workers' requires expression of integer type ('struct Incomplete' invalid)}}
+#pragma acc kernels num_workers(*SomeIncomplete)
+  while(1);
+
+  enum E{A} SomeE;
+
+#pragma acc kernels num_workers(SomeE)
+  while(1);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-num_workers-clause.cpp b/clang/test/SemaOpenACC/compute-construct-num_workers-clause.cpp
new file mode 100644
index 00000000000000..9449b77d092f4a
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-num_workers-clause.cpp
@@ -0,0 +1,133 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct NotConvertible{} NC;
+struct Incomplete *SomeIncomplete; // #INCOMPLETE
+enum E{} SomeE;
+enum class E2{} SomeE2;
+
+struct CorrectConvert {
+  operator int();
+} Convert;
+
+struct ExplicitConvertOnly {
+  explicit operator int() const; // #EXPL_CONV
+} Explicit;
+
+struct AmbiguousConvert{
+  operator int(); // #AMBIG_INT
+  operator short(); // #AMBIG_SHORT
+  operator float();
+} Ambiguous;
+
+void Test() {
+#pragma acc parallel num_workers(1)
+  while(1);
+#pragma acc kernels num_workers(1)
+  while(1);
+
+  // expected-error at +1{{OpenACC clause 'num_workers' requires expression of integer type ('struct NotConvertible' invalid}}
+#pragma acc parallel num_workers(NC)
+  while(1);
+
+  // expected-error at +2{{OpenACC integer expression has incomplete class type 'struct Incomplete'}}
+  // expected-note@#INCOMPLETE{{forward declaration of 'Incomplete'}}
+#pragma acc kernels num_workers(*SomeIncomplete)
+  while(1);
+
+#pragma acc parallel num_workers(SomeE)
+  while(1);
+
+  // expected-error at +1{{OpenACC clause 'num_workers' requires expression of integer type ('enum E2' invalid}}
+#pragma acc kernels num_workers(SomeE2)
+  while(1);
+
+#pragma acc parallel num_workers(Convert)
+  while(1);
+
+  // expected-error at +2{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+#pragma acc kernels num_workers(Explicit)
+  while(1);
+
+  // expected-error at +3{{multiple conversions from expression type 'struct AmbiguousConvert' to an integral type}}
+  // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+  // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc parallel num_workers(Ambiguous)
+  while(1);
+}
+
+struct HasInt {
+  using IntTy = int;
+  using ShortTy = short;
+  static constexpr int value = 1;
+  static constexpr AmbiguousConvert ACValue;
+  static constexpr ExplicitConvertOnly EXValue;
+
+  operator char();
+};
+
+template<typename T>
+void TestInst() {
+
+  // expected-error at +1{{no member named 'Invalid' in 'HasInt'}}
+#pragma acc parallel num_workers(HasInt::Invalid)
+  while (1);
+
+  // expected-error at +2{{no member named 'Invalid' in 'HasInt'}}
+  // expected-note@#INST{{in instantiation of function template specialization 'TestInst<HasInt>' requested here}}
+#pragma acc kernels num_workers(T::Invalid)
+  while (1);
+
+  // expected-error at +3{{multiple conversions from expression type 'const AmbiguousConvert' to an integral type}}
+  // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+  // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc parallel num_workers(HasInt::ACValue)
+  while (1);
+
+  // expected-error at +3{{multiple conversions from expression type 'const AmbiguousConvert' to an integral type}}
+  // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+  // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc kernels num_workers(T::ACValue)
+  while (1);
+
+  // expected-error at +2{{OpenACC integer expression type 'const ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+#pragma acc parallel num_workers(HasInt::EXValue)
+  while (1);
+
+  // expected-error at +2{{OpenACC integer expression type 'const ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+#pragma acc kernels num_workers(T::EXValue)
+  while (1);
+
+#pragma acc parallel num_workers(HasInt::value)
+  while (1);
+
+#pragma acc kernels num_workers(T::value)
+  while (1);
+
+#pragma acc parallel num_workers(HasInt::IntTy{})
+  while (1);
+
+#pragma acc kernels num_workers(typename T::ShortTy{})
+  while (1);
+
+#pragma acc parallel num_workers(HasInt::IntTy{})
+  while (1);
+
+#pragma acc kernels num_workers(typename T::ShortTy{})
+  while (1);
+
+  HasInt HI{};
+  T MyT{};
+
+#pragma acc parallel num_workers(HI)
+  while (1);
+
+#pragma acc kernels num_workers(MyT)
+  while (1);
+}
+
+void Inst() {
+  TestInst<HasInt>(); // #INST
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 2ef599d2cd26fa..011bec32bab319 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2795,6 +2795,10 @@ void OpenACCClauseEnqueue::VisitSelfClause(const OpenACCSelfClause &C) {
   if (C.hasConditionExpr())
     Visitor.AddStmt(C.getConditionExpr());
 }
+void OpenACCClauseEnqueue::VisitNumWorkersClause(
+    const OpenACCNumWorkersClause &C) {
+  Visitor.AddStmt(C.getIntExpr());
+}
 } // namespace
 
 void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {



More information about the cfe-commits mailing list