[clang] [OpenACC] Implement 'num_gangs' sema for compute constructs (PR #89460)

Erich Keane via cfe-commits cfe-commits at lists.llvm.org
Mon Apr 22 08:57:22 PDT 2024


https://github.com/erichkeane updated https://github.com/llvm/llvm-project/pull/89460

>From aa7844c73b72172707ec81234c8e9d5370bbb772 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Fri, 19 Apr 2024 10:20:40 -0700
Subject: [PATCH 1/2] [OpenACC] Implement 'num_gangs' sema for compute
 constructs

num_gangs takes an 'int-expr-list', for 'parallel', and an 'int-expr' for
'kernels'.  This patch changes the parsing to always parse it as an
'int-expr-list', then correct the expression count during Sema.  It also
implements the rest of the semantic analysis changes for this clause.
---
 clang/include/clang/AST/OpenACCClause.h       |  87 ++++++++--
 .../clang/Basic/DiagnosticSemaKinds.td        |   5 +
 clang/include/clang/Basic/OpenACCClauses.def  |   1 +
 clang/include/clang/Parse/Parser.h            |  16 +-
 clang/include/clang/Sema/SemaOpenACC.h        |  16 +-
 clang/lib/AST/OpenACCClause.cpp               |  16 ++
 clang/lib/AST/StmtProfile.cpp                 |   6 +
 clang/lib/AST/TextNodeDumper.cpp              |   1 +
 clang/lib/Parse/ParseOpenACC.cpp              |  87 ++++++++--
 clang/lib/Sema/SemaOpenACC.cpp                |  35 ++++
 clang/lib/Sema/TreeTransform.h                |  26 +++
 clang/lib/Serialization/ASTReader.cpp         |  10 +-
 clang/lib/Serialization/ASTWriter.cpp         |   9 +-
 clang/test/ParserOpenACC/parse-clauses.c      |   4 -
 .../compute-construct-intexpr-clause-ast.cpp  |  79 +++++++++
 .../compute-construct-num_gangs-clause.c      |  54 ++++++
 .../compute-construct-num_gangs-clause.cpp    | 160 ++++++++++++++++++
 clang/tools/libclang/CIndex.cpp               |   4 +
 18 files changed, 572 insertions(+), 44 deletions(-)
 create mode 100644 clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
 create mode 100644 clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp

diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 8b6d3221aa066b..277a351c49fcb8 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -156,33 +156,88 @@ class OpenACCSelfClause : public OpenACCClauseWithCondition {
                                    Expr *ConditionExpr, SourceLocation EndLoc);
 };
 
-/// Represents oen of a handful of classes that have a single integer
+/// Represents a clause that has one or more IntExprs.  It does not own the
+/// IntExprs, but provides 'children' and other accessors.
+class OpenACCClauseWithIntExprs : public OpenACCClauseWithParams {
+  MutableArrayRef<Expr *> IntExprs;
+
+protected:
+  OpenACCClauseWithIntExprs(OpenACCClauseKind K, SourceLocation BeginLoc,
+                            SourceLocation LParenLoc, SourceLocation EndLoc)
+      : OpenACCClauseWithParams(K, BeginLoc, LParenLoc, EndLoc) {}
+
+  /// Used only for initialization, the leaf class can initialize this to
+  /// trailing storage.
+  void setIntExprs(MutableArrayRef<Expr *> NewIntExprs) {
+    assert(IntExprs.empty() && "Cannot change IntExprs list");
+    IntExprs = NewIntExprs;
+  }
+
+  /// 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());
+  }
+};
+
+class OpenACCNumGangsClause final
+    : public OpenACCClauseWithIntExprs,
+      public llvm::TrailingObjects<OpenACCNumGangsClause, Expr *> {
+
+  OpenACCNumGangsClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                        ArrayRef<Expr *> IntExprs, SourceLocation EndLoc)
+      : OpenACCClauseWithIntExprs(OpenACCClauseKind::NumGangs, BeginLoc,
+                                  LParenLoc, EndLoc) {
+    std::uninitialized_copy(IntExprs.begin(), IntExprs.end(),
+                            getTrailingObjects<Expr *>());
+    setIntExprs(MutableArrayRef(getTrailingObjects<Expr *>(), IntExprs.size()));
+  }
+
+public:
+  static OpenACCNumGangsClause *
+  Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+         ArrayRef<Expr *> IntExprs, SourceLocation EndLoc);
+
+  llvm::ArrayRef<Expr *> getIntExprs() {
+    return OpenACCClauseWithIntExprs::getIntExprs();
+  }
+
+  llvm::ArrayRef<Expr *> getIntExprs() const {
+    return OpenACCClauseWithIntExprs::getIntExprs();
+  }
+};
+
+/// Represents one of a handful of clauses that have a single integer
 /// expression.
-class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithParams {
+class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithIntExprs {
   Expr *IntExpr;
 
 protected:
   OpenACCClauseWithSingleIntExpr(OpenACCClauseKind K, SourceLocation BeginLoc,
                                  SourceLocation LParenLoc, Expr *IntExpr,
                                  SourceLocation EndLoc)
-      : OpenACCClauseWithParams(K, BeginLoc, LParenLoc, EndLoc),
-        IntExpr(IntExpr) {}
+      : OpenACCClauseWithIntExprs(K, BeginLoc, LParenLoc, EndLoc),
+        IntExpr(IntExpr) {
+    setIntExprs(MutableArrayRef<Expr *>{&this->IntExpr, 1});
+  }
 
 public:
-  bool hasIntExpr() const { return IntExpr; }
-  const Expr *getIntExpr() const { return IntExpr; }
-
-  Expr *getIntExpr() { return IntExpr; };
-
-  child_range children() {
-    return child_range(reinterpret_cast<Stmt **>(&IntExpr),
-                       reinterpret_cast<Stmt **>(&IntExpr + 1));
+  bool hasIntExpr() const { return !getIntExprs().empty(); }
+  const Expr *getIntExpr() const {
+    return hasIntExpr() ? getIntExprs()[0] : nullptr;
   }
 
-  const_child_range children() const {
-    return const_child_range(reinterpret_cast<Stmt *const *>(&IntExpr),
-                             reinterpret_cast<Stmt *const *>(&IntExpr + 1));
-  }
+  Expr *getIntExpr() { return hasIntExpr() ? getIntExprs()[0] : nullptr; };
 };
 
 class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr {
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 1bbe76ff6bd2ac..a95424862e63f4 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12293,4 +12293,9 @@ 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">;
+def err_acc_num_gangs_num_args
+    : Error<"%select{no|too many}0 integer expression arguments provided to "
+            "OpenACC 'num_gangs' "
+            "%select{|clause: '%1' directive expects maximum of %2, %3 were "
+            "provided}0">;
 } // end of sema component.
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 520e068f4ffd40..dd5792e7ca8c39 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -18,6 +18,7 @@
 VISIT_CLAUSE(Default)
 VISIT_CLAUSE(If)
 VISIT_CLAUSE(Self)
+VISIT_CLAUSE(NumGangs)
 VISIT_CLAUSE(NumWorkers)
 VISIT_CLAUSE(VectorLength)
 
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 72b2f958a5e622..7e95a52e6f34d5 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3644,10 +3644,22 @@ class Parser : public CodeCompletionHandler {
   /// Parses the clause of the 'bind' argument, which can be a string literal or
   /// an ID expression.
   ExprResult ParseOpenACCBindClauseArgument();
+
+  /// A type to represent the state of parsing after an attempt to parse an
+  /// OpenACC int-expr. This is useful to determine whether an int-expr list can
+  /// continue parsing after a failed int-expr.
+  using OpenACCIntExprParseResult =
+      std::pair<ExprResult, OpenACCParseCanContinue>;
   /// Parses the clause kind of 'int-expr', which can be any integral
   /// expression.
-  ExprResult ParseOpenACCIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
-                                 SourceLocation Loc);
+  OpenACCIntExprParseResult ParseOpenACCIntExpr(OpenACCDirectiveKind DK,
+                                                OpenACCClauseKind CK,
+                                                SourceLocation Loc);
+  /// Parses the argument list for 'num_gangs', which allows up to 3
+  /// 'int-expr's.
+  bool ParseOpenACCIntExprList(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+                               SourceLocation Loc,
+                               llvm::SmallVector<Expr *> &IntExprs);
   /// Parses the 'device-type-list', which is a list of identifiers.
   bool ParseOpenACCDeviceTypeList();
   /// Parses the 'async-argument', which is an integral value with two
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 023722049732af..ea28617f79b81b 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -93,14 +93,16 @@ class SemaOpenACC : public SemaBase {
     }
 
     unsigned getNumIntExprs() const {
-      assert((ClauseKind == OpenACCClauseKind::NumWorkers ||
+      assert((ClauseKind == OpenACCClauseKind::NumGangs ||
+              ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
       return std::get<IntExprDetails>(Details).IntExprs.size();
     }
 
     ArrayRef<Expr *> getIntExprs() {
-      assert((ClauseKind == OpenACCClauseKind::NumWorkers ||
+      assert((ClauseKind == OpenACCClauseKind::NumGangs ||
+              ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
       return std::get<IntExprDetails>(Details).IntExprs;
@@ -134,11 +136,19 @@ class SemaOpenACC : public SemaBase {
     }
 
     void setIntExprDetails(ArrayRef<Expr *> IntExprs) {
-      assert((ClauseKind == OpenACCClauseKind::NumWorkers ||
+      assert((ClauseKind == OpenACCClauseKind::NumGangs ||
+              ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
       Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}};
     }
+    void setIntExprDetails(llvm::SmallVector<Expr *> &&IntExprs) {
+      assert((ClauseKind == OpenACCClauseKind::NumGangs ||
+              ClauseKind == OpenACCClauseKind::NumWorkers ||
+              ClauseKind == OpenACCClauseKind::VectorLength) &&
+             "Parsed clause kind does not have a int exprs");
+      Details = IntExprDetails{IntExprs};
+    }
   };
 
   SemaOpenACC(Sema &S);
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 75334223e073c3..6cd5b28802187d 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -124,6 +124,16 @@ OpenACCVectorLengthClause::Create(const ASTContext &C, SourceLocation BeginLoc,
       OpenACCVectorLengthClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
 }
 
+OpenACCNumGangsClause *OpenACCNumGangsClause::Create(const ASTContext &C,
+                                                     SourceLocation BeginLoc,
+                                                     SourceLocation LParenLoc,
+                                                     ArrayRef<Expr *> IntExprs,
+                                                     SourceLocation EndLoc) {
+  void *Mem = C.Allocate(
+      OpenACCNumGangsClause::totalSizeToAlloc<Expr *>(IntExprs.size()));
+  return new (Mem) OpenACCNumGangsClause(BeginLoc, LParenLoc, IntExprs, EndLoc);
+}
+
 //===----------------------------------------------------------------------===//
 //  OpenACC clauses printing methods
 //===----------------------------------------------------------------------===//
@@ -141,6 +151,12 @@ void OpenACCClausePrinter::VisitSelfClause(const OpenACCSelfClause &C) {
     OS << "(" << CondExpr << ")";
 }
 
+void OpenACCClausePrinter::VisitNumGangsClause(const OpenACCNumGangsClause &C) {
+  OS << "num_gangs(";
+  llvm::interleaveComma(C.getIntExprs(), OS);
+  OS << ")";
+}
+
 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 8138ae3a244a8b..c81724f84dd9ce 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2497,6 +2497,12 @@ void OpenACCClauseProfiler::VisitSelfClause(const OpenACCSelfClause &Clause) {
     Profiler.VisitStmt(Clause.getConditionExpr());
 }
 
+void OpenACCClauseProfiler::VisitNumGangsClause(
+    const OpenACCNumGangsClause &Clause) {
+  for (auto *E : Clause.getIntExprs())
+    Profiler.VisitStmt(E);
+}
+
 void OpenACCClauseProfiler::VisitNumWorkersClause(
     const OpenACCNumWorkersClause &Clause) {
   assert(Clause.hasIntExpr() && "num_workers clause requires a valid int expr");
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index e5a8b285715b30..8f0a9a9b0ed0bc 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::NumGangs:
     case OpenACCClauseKind::NumWorkers:
     case OpenACCClauseKind::VectorLength:
       // The condition expression will be printed as a part of the 'children',
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 757417f75c9636..9b7cd1668f7c17 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -632,16 +632,54 @@ Parser::ParseOpenACCClauseList(OpenACCDirectiveKind DirKind) {
   return Clauses;
 }
 
-ExprResult Parser::ParseOpenACCIntExpr(OpenACCDirectiveKind DK,
-                                       OpenACCClauseKind CK,
-                                       SourceLocation Loc) {
-  ExprResult ER =
-      getActions().CorrectDelayedTyposInExpr(ParseAssignmentExpression());
+Parser::OpenACCIntExprParseResult
+Parser::ParseOpenACCIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+                            SourceLocation Loc) {
+  ExprResult ER = ParseAssignmentExpression();
 
+  // If the actual parsing failed, we don't know the state of the parse, so
+  // don't try to continue.
   if (!ER.isUsable())
-    return ER;
+    return {ER, OpenACCParseCanContinue::Cannot};
+
+  // Parsing can continue after the initial assignment expression parsing, so
+  // even if there was a typo, we can continue.
+  ER = getActions().CorrectDelayedTyposInExpr(ER);
+  if (!ER.isUsable())
+    return {ER, OpenACCParseCanContinue::Can};
+
+  return {getActions().OpenACC().ActOnIntExpr(DK, CK, Loc, ER.get()),
+          OpenACCParseCanContinue::Can};
+}
+
+bool Parser::ParseOpenACCIntExprList(OpenACCDirectiveKind DK,
+                                     OpenACCClauseKind CK, SourceLocation Loc,
+                                     llvm::SmallVector<Expr *> &IntExprs) {
+  OpenACCIntExprParseResult CurResult = ParseOpenACCIntExpr(DK, CK, Loc);
+
+  if (!CurResult.first.isUsable() &&
+      CurResult.second == OpenACCParseCanContinue::Cannot) {
+    SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
+              Parser::StopBeforeMatch);
+    return true;
+  }
+
+  IntExprs.push_back(CurResult.first.get());
+
+  while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
+    ExpectAndConsume(tok::comma);
+
+    CurResult = ParseOpenACCIntExpr(DK, CK, Loc);
 
-  return getActions().OpenACC().ActOnIntExpr(DK, CK, Loc, ER.get());
+    if (!CurResult.first.isUsable() &&
+        CurResult.second == OpenACCParseCanContinue::Cannot) {
+      SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
+                Parser::StopBeforeMatch);
+      return true;
+    }
+    IntExprs.push_back(CurResult.first.get());
+  }
+  return false;
 }
 
 bool Parser::ParseOpenACCClauseVarList(OpenACCClauseKind Kind) {
@@ -761,7 +799,7 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
     ConsumeToken();
     return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
                                OpenACCClauseKind::Gang, GangLoc)
-        .isInvalid();
+        .first.isInvalid();
   }
 
   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Num, getCurToken()) &&
@@ -773,7 +811,7 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
   // This is just the 'num' case where 'num' is optional.
   return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
                              OpenACCClauseKind::Gang, GangLoc)
-      .isInvalid();
+      .first.isInvalid();
 }
 
 bool Parser::ParseOpenACCGangArgList(SourceLocation GangLoc) {
@@ -946,13 +984,25 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
       }
       break;
     }
-    case OpenACCClauseKind::NumGangs:
+    case OpenACCClauseKind::NumGangs: {
+      llvm::SmallVector<Expr *> IntExprs;
+
+      if (ParseOpenACCIntExprList(OpenACCDirectiveKind::Invalid,
+                                  OpenACCClauseKind::NumGangs, ClauseLoc,
+                                  IntExprs)) {
+        Parens.skipToEnd();
+        return OpenACCCanContinue();
+      }
+      ParsedClause.setIntExprDetails(std::move(IntExprs));
+      break;
+    }
     case OpenACCClauseKind::NumWorkers:
     case OpenACCClauseKind::DeviceNum:
     case OpenACCClauseKind::DefaultAsync:
     case OpenACCClauseKind::VectorLength: {
       ExprResult IntExpr = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
-                                               ClauseKind, ClauseLoc);
+                                               ClauseKind, ClauseLoc)
+                               .first;
       if (IntExpr.isInvalid()) {
         Parens.skipToEnd();
         return OpenACCCanContinue();
@@ -1017,7 +1067,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
                                                : OpenACCSpecialTokenKind::Num,
                                            ClauseKind);
         ExprResult IntExpr = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
-                                                 ClauseKind, ClauseLoc);
+                                                 ClauseKind, ClauseLoc)
+                                 .first;
         if (IntExpr.isInvalid()) {
           Parens.skipToEnd();
           return OpenACCCanContinue();
@@ -1081,11 +1132,13 @@ bool Parser::ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective) {
     // Consume colon.
     ConsumeToken();
 
-    ExprResult IntExpr = ParseOpenACCIntExpr(
-        IsDirective ? OpenACCDirectiveKind::Wait
-                    : OpenACCDirectiveKind::Invalid,
-        IsDirective ? OpenACCClauseKind::Invalid : OpenACCClauseKind::Wait,
-        Loc);
+    ExprResult IntExpr =
+        ParseOpenACCIntExpr(IsDirective ? OpenACCDirectiveKind::Wait
+                                        : OpenACCDirectiveKind::Invalid,
+                            IsDirective ? OpenACCClauseKind::Invalid
+                                        : OpenACCClauseKind::Wait,
+                            Loc)
+            .first;
     if (IntExpr.isInvalid())
       return true;
 
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 190739fa02e932..ba69e71e30a181 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -91,6 +91,7 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
     default:
       return false;
     }
+  case OpenACCClauseKind::NumGangs:
   case OpenACCClauseKind::NumWorkers:
   case OpenACCClauseKind::VectorLength:
     switch (DirectiveKind) {
@@ -230,6 +231,40 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
         getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
         Clause.getConditionExpr(), Clause.getEndLoc());
   }
+  case OpenACCClauseKind::NumGangs: {
+    // 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;
+
+    if (Clause.getIntExprs().empty())
+      Diag(Clause.getBeginLoc(), diag::err_acc_num_gangs_num_args)
+          << /*NoArgs=*/0;
+
+    unsigned MaxArgs =
+        (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel ||
+         Clause.getDirectiveKind() == OpenACCDirectiveKind::ParallelLoop)
+            ? 3
+            : 1;
+    if (Clause.getIntExprs().size() > MaxArgs)
+      Diag(Clause.getBeginLoc(), diag::err_acc_num_gangs_num_args)
+          << /*NoArgs=*/1 << Clause.getDirectiveKind() << MaxArgs
+          << Clause.getIntExprs().size();
+
+    // Create the AST node for the clause even if the number of expressions is
+    // incorrect.
+    return OpenACCNumGangsClause::Create(
+        getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
+        Clause.getIntExprs(), Clause.getEndLoc());
+    break;
+  }
   case OpenACCClauseKind::NumWorkers: {
     // Restrictions only properly implemented on 'compute' constructs, and
     // 'compute' constructs are the only construct that can do anything with
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index ade33ec65038fd..10ce226f1298dd 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11159,6 +11159,32 @@ void OpenACCClauseTransform<Derived>::VisitSelfClause(
       ParsedClause.getEndLoc());
 }
 
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitNumGangsClause(
+    const OpenACCNumGangsClause &C) {
+  llvm::SmallVector<Expr *> InstantiatedIntExprs;
+
+  for (Expr *CurIntExpr : C.getIntExprs()) {
+    ExprResult Res = Self.TransformExpr(CurIntExpr);
+
+    if (!Res.isUsable())
+      return;
+
+    Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid,
+                                                C.getClauseKind(),
+                                                C.getBeginLoc(), Res.get());
+    if (!Res.isUsable())
+      return;
+
+    InstantiatedIntExprs.push_back(Res.get());
+  }
+
+  ParsedClause.setIntExprDetails(InstantiatedIntExprs);
+  NewClause = OpenACCNumGangsClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getIntExprs(),
+      ParsedClause.getEndLoc());
+}
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitNumWorkersClause(
     const OpenACCNumWorkersClause &C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 44e23919ea18e0..d64925676df7b1 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11786,6 +11786,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCSelfClause::Create(getContext(), BeginLoc, LParenLoc,
                                      CondExpr, EndLoc);
   }
+  case OpenACCClauseKind::NumGangs: {
+    SourceLocation LParenLoc = readSourceLocation();
+    unsigned NumClauses = readInt();
+    llvm::SmallVector<Expr *> IntExprs;
+    for (unsigned I = 0; I < NumClauses; ++I)
+      IntExprs.push_back(readSubExpr());
+    return OpenACCNumGangsClause::Create(getContext(), BeginLoc, LParenLoc,
+                                         IntExprs, EndLoc);
+  }
   case OpenACCClauseKind::NumWorkers: {
     SourceLocation LParenLoc = readSourceLocation();
     Expr *IntExpr = readSubExpr();
@@ -11826,7 +11835,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   case OpenACCClauseKind::Reduction:
   case OpenACCClauseKind::Collapse:
   case OpenACCClauseKind::Bind:
-  case OpenACCClauseKind::NumGangs:
   case OpenACCClauseKind::DeviceNum:
   case OpenACCClauseKind::DefaultAsync:
   case OpenACCClauseKind::DeviceType:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 8a4b36207c4734..018b854652a46e 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -7657,6 +7657,14 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
       AddStmt(const_cast<Expr *>(SC->getConditionExpr()));
     return;
   }
+  case OpenACCClauseKind::NumGangs: {
+    const auto *NGC = cast<OpenACCNumGangsClause>(C);
+    writeSourceLocation(NGC->getLParenLoc());
+    writeUInt32(NGC->getIntExprs().size());
+    for (Expr *E : NGC->getIntExprs())
+      AddStmt(E);
+    return;
+  }
   case OpenACCClauseKind::NumWorkers: {
     const auto *NWC = cast<OpenACCNumWorkersClause>(C);
     writeSourceLocation(NWC->getLParenLoc());
@@ -7697,7 +7705,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::Reduction:
   case OpenACCClauseKind::Collapse:
   case OpenACCClauseKind::Bind:
-  case OpenACCClauseKind::NumGangs:
   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 ddf40f71701eda..799f22b8c120e5 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -911,16 +911,12 @@ void IntExprParsing() {
 #pragma acc parallel num_gangs(invalid)
   {}
 
-  // expected-error at +2{{expected ')'}}
-  // expected-note at +1{{to match this '('}}
 #pragma acc parallel num_gangs(5, 4)
   {}
 
-  // expected-warning at +1{{OpenACC clause 'num_gangs' not yet implemented, clause ignored}}
 #pragma acc parallel num_gangs(5)
   {}
 
-  // expected-warning at +1{{OpenACC clause 'num_gangs' not yet implemented, clause ignored}}
 #pragma acc parallel num_gangs(returns_int())
   {}
 
diff --git a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
index 889025c26818d8..5a4c9f05ee089e 100644
--- a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
@@ -88,6 +88,34 @@ void NormalUses() {
   // CHECK-NEXT: WhileStmt
   // CHECK-NEXT: CXXBoolLiteralExpr
   // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_gangs(some_int(), some_long(), some_short())
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_gangs clause
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
+  // CHECK-NEXT: CallExpr{{.*}}'long'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'long (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'long ()' lvalue Function{{.*}} 'some_long' 'long ()'
+  // 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 kernels num_gangs(some_int())
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_gangs 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
 }
 
 template<typename T, typename U>
@@ -187,6 +215,31 @@ void TemplUses(T t, U u) {
   // CHECK-NEXT: CXXBoolLiteralExpr
   // CHECK-NEXT: CompoundStmt
 
+#pragma acc kernels num_gangs(u)
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_gangs clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_gangs(u, U::value)
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: num_gangs clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}}EndMarker
+  int EndMarker;
+
   // Check the instantiated versions of the above.
   // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation
   // CHECK-NEXT: TemplateArgument type 'CorrectConvert'
@@ -288,6 +341,32 @@ void TemplUses(T t, U u) {
   // CHECK-NEXT: WhileStmt
   // CHECK-NEXT: CXXBoolLiteralExpr
   // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: num_gangs 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_gangs 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: 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: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}}EndMarker
 }
 
 struct HasInt {
diff --git a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
new file mode 100644
index 00000000000000..cdc6847b47f948
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
@@ -0,0 +1,54 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+short getS();
+void Test() {
+#pragma acc kernels num_gangs(1)
+  while(1);
+
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1)
+  while(1);
+
+#pragma acc parallel num_gangs(1)
+  while(1);
+
+  // expected-error at +2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels' directive}}
+  // expected-note at +1{{previous clause is here}}
+#pragma acc kernels num_gangs(1) num_gangs(2)
+  while(1);
+
+  // expected-error at +2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel' directive}}
+  // expected-note at +1{{previous clause is here}}
+#pragma acc parallel num_gangs(1) num_gangs(2)
+  while(1);
+
+  // expected-error at +1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels' directive expects maximum of 1, 2 were provided}}
+#pragma acc kernels num_gangs(1, getS())
+  while(1);
+
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1, getS())
+  while(1);
+#pragma acc parallel num_gangs(1, getS())
+  while(1);
+
+  struct NotConvertible{} NC;
+  // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(NC)
+  while(1);
+
+  // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(1, NC)
+  while(1);
+
+  // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(NC, 1)
+  while(1);
+
+#pragma acc parallel num_gangs(getS(), 1, getS())
+  while(1);
+
+  // expected-error at +1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'parallel' directive expects maximum of 3, 4 were provided}}
+#pragma acc parallel num_gangs(getS(), 1, getS(), 1)
+  while(1);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp
new file mode 100644
index 00000000000000..ec3df87a065572
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp
@@ -0,0 +1,160 @@
+// 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;
+
+short some_short();
+int some_int();
+long some_long();
+
+void Test() {
+#pragma acc kernels num_gangs(1)
+  while(1);
+
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1)
+  while(1);
+
+#pragma acc parallel num_gangs(1)
+  while(1);
+
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(some_short(), some_int(), some_long())
+  while(1);
+
+#pragma acc parallel num_gangs(some_short(), some_int(), some_long())
+  while(1);
+
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(some_short(), some_int(), some_long(), SomeE)
+  while(1);
+
+  // expected-error at +1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'parallel' directive expects maximum of 3, 4 were provided}}
+#pragma acc parallel num_gangs(some_short(), some_int(), some_long(), SomeE)
+  while(1);
+
+  // expected-error at +1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels' directive expects maximum of 1, 2 were provided}}
+#pragma acc kernels num_gangs(1, 2)
+  while(1);
+
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1, 2)
+  while(1);
+
+#pragma acc parallel num_gangs(1, 2)
+  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_gangs(Ambiguous)
+  while(1);
+
+  // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(NC, SomeE)
+  while(1);
+
+  // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(SomeE, NC)
+  while(1);
+
+  // expected-error at +3{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+  // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(Explicit, NC)
+  while(1);
+
+  // expected-error at +4{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+  // expected-error at +2{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(Explicit, NC)
+  while(1);
+
+  // expected-error at +6{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+  // expected-error at +4{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+  // 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_gangs(Explicit, NC, Ambiguous)
+  while(1);
+
+  // expected-error at +7{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+  // expected-error at +5{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+  // expected-error at +4{{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'}}
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(Explicit, NC, Ambiguous)
+  while(1);
+  // TODO
+}
+
+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 +2{{no member named 'Invalid' in 'HasInt'}}
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(HasInt::Invalid)
+  while(1);
+
+  // expected-error at +2{{no member named 'Invalid' in 'HasInt'}}
+  // expected-note@#INST{{in instantiation of function template specialization}}
+#pragma acc parallel num_gangs(T::Invalid)
+  while(1);
+
+  // expected-error at +1{{no member named 'Invalid' in 'HasInt'}}
+#pragma acc parallel num_gangs(1, HasInt::Invalid)
+  while(1);
+
+  // expected-error at +1{{no member named 'Invalid' in 'HasInt'}}
+#pragma acc parallel num_gangs(T::Invalid, 1)
+  while(1);
+
+  // expected-error at +2{{no member named 'Invalid' in 'HasInt'}}
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1, HasInt::Invalid)
+  while(1);
+
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(T::Invalid, 1)
+  while(1);
+
+#pragma acc parallel num_gangs(T::value, typename T::IntTy{})
+  while(1);
+
+  // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(T::value, typename T::IntTy{})
+  while(1);
+}
+
+void Inst() {
+  TestInst<HasInt>(); // #INST
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index cbc1d85bb33dfc..74163f30e19b1d 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2803,6 +2803,10 @@ void OpenACCClauseEnqueue::VisitVectorLengthClause(
     const OpenACCVectorLengthClause &C) {
   Visitor.AddStmt(C.getIntExpr());
 }
+void OpenACCClauseEnqueue::VisitNumGangsClause(const OpenACCNumGangsClause &C) {
+  for (Expr *IE : C.getIntExprs())
+    Visitor.AddStmt(IE);
+}
 } // namespace
 
 void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {

>From 9e1dcd21c965df63c526d09a346f376350770105 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Mon, 22 Apr 2024 08:43:31 -0700
Subject: [PATCH 2/2] SmallVector to SmallVectorImpl

---
 clang/include/clang/Parse/Parser.h | 2 +-
 clang/lib/Parse/ParseOpenACC.cpp   | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 7e95a52e6f34d5..d3bb04ff7a2c6d 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3659,7 +3659,7 @@ class Parser : public CodeCompletionHandler {
   /// 'int-expr's.
   bool ParseOpenACCIntExprList(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
                                SourceLocation Loc,
-                               llvm::SmallVector<Expr *> &IntExprs);
+                               llvm::SmallVectorImpl<Expr *> &IntExprs);
   /// Parses the 'device-type-list', which is a list of identifiers.
   bool ParseOpenACCDeviceTypeList();
   /// Parses the 'async-argument', which is an integral value with two
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 9b7cd1668f7c17..8a18fca8064ee1 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -654,7 +654,7 @@ Parser::ParseOpenACCIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
 
 bool Parser::ParseOpenACCIntExprList(OpenACCDirectiveKind DK,
                                      OpenACCClauseKind CK, SourceLocation Loc,
-                                     llvm::SmallVector<Expr *> &IntExprs) {
+                                     llvm::SmallVectorImpl<Expr *> &IntExprs) {
   OpenACCIntExprParseResult CurResult = ParseOpenACCIntExpr(DK, CK, Loc);
 
   if (!CurResult.first.isUsable() &&



More information about the cfe-commits mailing list