[clang] 5b25c31 - [OpenACC] Implement loop 'gang' clause. (#112006)

via cfe-commits cfe-commits at lists.llvm.org
Fri Oct 11 09:05:23 PDT 2024


Author: Erich Keane
Date: 2024-10-11T09:05:19-07:00
New Revision: 5b25c31351ad1b10a3819411379b3258869c1e1b

URL: https://github.com/llvm/llvm-project/commit/5b25c31351ad1b10a3819411379b3258869c1e1b
DIFF: https://github.com/llvm/llvm-project/commit/5b25c31351ad1b10a3819411379b3258869c1e1b.diff

LOG: [OpenACC] Implement loop 'gang' clause. (#112006)

The 'gang' clause is used to specify parallel execution of loops, thus
has some complicated rules depending on the 'loop's associated compute
construct. This patch implements all of those.

Added: 
    clang/test/SemaOpenACC/loop-construct-gang-ast.cpp
    clang/test/SemaOpenACC/loop-construct-gang-clause.cpp

Modified: 
    clang/include/clang/AST/OpenACCClause.h
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/include/clang/Basic/OpenACCClauses.def
    clang/include/clang/Basic/OpenACCKinds.h
    clang/include/clang/Parse/Parser.h
    clang/include/clang/Sema/SemaOpenACC.h
    clang/lib/AST/OpenACCClause.cpp
    clang/lib/AST/StmtProfile.cpp
    clang/lib/AST/TextNodeDumper.cpp
    clang/lib/Parse/ParseOpenACC.cpp
    clang/lib/Sema/SemaOpenACC.cpp
    clang/lib/Sema/TreeTransform.h
    clang/lib/Serialization/ASTReader.cpp
    clang/lib/Serialization/ASTWriter.cpp
    clang/test/AST/ast-print-openacc-loop-construct.cpp
    clang/test/ParserOpenACC/parse-clauses.c
    clang/test/SemaOpenACC/compute-construct-device_type-clause.c
    clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
    clang/test/SemaOpenACC/loop-construct-device_type-clause.c
    clang/tools/libclang/CIndex.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index b500acc768e55a..f3a09eb651458d 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -119,32 +119,6 @@ class OpenACCSeqClause : 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 OpenACCGangClause : public OpenACCClause {
-protected:
-  OpenACCGangClause(SourceLocation BeginLoc, SourceLocation EndLoc)
-      : OpenACCClause(OpenACCClauseKind::Gang, BeginLoc, EndLoc) {
-    llvm_unreachable("Not yet implemented");
-  }
-
-public:
-  static bool classof(const OpenACCClause *C) {
-    return C->getClauseKind() == OpenACCClauseKind::Gang;
-  }
-
-  static OpenACCGangClause *
-  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());
-  }
-};
-
 // 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.
@@ -177,7 +151,7 @@ class OpenACCVectorClause : public OpenACCClause {
 class OpenACCWorkerClause : public OpenACCClause {
 protected:
   OpenACCWorkerClause(SourceLocation BeginLoc, SourceLocation EndLoc)
-      : OpenACCClause(OpenACCClauseKind::Gang, BeginLoc, EndLoc) {
+      : OpenACCClause(OpenACCClauseKind::Worker, BeginLoc, EndLoc) {
     llvm_unreachable("Not yet implemented");
   }
 
@@ -535,6 +509,38 @@ class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithExprs {
   Expr *getIntExpr() { return hasIntExpr() ? getExprs()[0] : nullptr; };
 };
 
+class OpenACCGangClause final
+    : public OpenACCClauseWithExprs,
+      public llvm::TrailingObjects<OpenACCGangClause, Expr *, OpenACCGangKind> {
+protected:
+  OpenACCGangClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                    ArrayRef<OpenACCGangKind> GangKinds,
+                    ArrayRef<Expr *> IntExprs, SourceLocation EndLoc);
+
+  OpenACCGangKind getGangKind(unsigned I) const {
+    return getTrailingObjects<OpenACCGangKind>()[I];
+  }
+
+public:
+  static bool classof(const OpenACCClause *C) {
+    return C->getClauseKind() == OpenACCClauseKind::Gang;
+  }
+
+  size_t numTrailingObjects(OverloadToken<Expr *>) const {
+    return getNumExprs();
+  }
+
+  unsigned getNumExprs() const { return getExprs().size(); }
+  std::pair<OpenACCGangKind, const Expr *> getExpr(unsigned I) const {
+    return {getGangKind(I), getExprs()[I]};
+  }
+
+  static OpenACCGangClause *
+  Create(const ASTContext &Ctx, SourceLocation BeginLoc,
+         SourceLocation LParenLoc, ArrayRef<OpenACCGangKind> GangKinds,
+         ArrayRef<Expr *> IntExprs, 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 41cdd09e971651..3c62a017005e59 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12576,6 +12576,7 @@ def err_acc_duplicate_clause_disallowed
     : Error<"OpenACC '%1' clause cannot appear more than once on a '%0' "
             "directive">;
 def note_acc_previous_clause_here : Note<"previous clause is here">;
+def note_acc_previous_expr_here : Note<"previous expression is here">;
 def err_acc_branch_in_out_compute_construct
     : Error<"invalid %select{branch|return|throw}0 %select{out of|into}1 "
             "OpenACC Compute Construct">;
@@ -12682,6 +12683,24 @@ def err_acc_insufficient_loops
 def err_acc_intervening_code
     : Error<"inner loops must be tightly nested inside a '%0' clause on "
             "a 'loop' construct">;
+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 "
+            "'parallel' compute construct|associated with a 'kernels' compute "
+            "construct|associated with a 'serial' compute construct}1">;
+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' "
+            "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">;
 
 // 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 a380e5ae69c418..2a098de31eb618 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -42,6 +42,7 @@ VISIT_CLAUSE(DevicePtr)
 VISIT_CLAUSE(DeviceType)
 CLAUSE_ALIAS(DType, DeviceType, false)
 VISIT_CLAUSE(FirstPrivate)
+VISIT_CLAUSE(Gang)
 VISIT_CLAUSE(If)
 VISIT_CLAUSE(Independent)
 VISIT_CLAUSE(NoCreate)

diff  --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h
index c4dfe3bedc13a7..3f48ebca708a42 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/clang/include/clang/Basic/OpenACCKinds.h
@@ -550,6 +550,35 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
                                      OpenACCReductionOperator Op) {
   return printOpenACCReductionOperator(Out, Op);
 }
+
+enum class OpenACCGangKind : uint8_t {
+  /// num:
+  Num,
+  /// dim:
+  Dim,
+  /// static:
+  Static
+};
+
+template <typename StreamTy>
+inline StreamTy &printOpenACCGangKind(StreamTy &Out, OpenACCGangKind GK) {
+  switch (GK) {
+  case OpenACCGangKind::Num:
+    return Out << "num";
+  case OpenACCGangKind::Dim:
+    return Out << "dim";
+  case OpenACCGangKind::Static:
+    return Out << "static";
+  }
+}
+inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
+                                             OpenACCGangKind Op) {
+  return printOpenACCGangKind(Out, Op);
+}
+inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
+                                     OpenACCGangKind Op) {
+  return printOpenACCGangKind(Out, Op);
+}
 } // namespace clang
 
 #endif // LLVM_CLANG_BASIC_OPENACCKINDS_H

diff  --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index dbcb545058a026..045ee754a242b3 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3797,9 +3797,15 @@ class Parser : public CodeCompletionHandler {
   bool ParseOpenACCSizeExprList(OpenACCClauseKind CK,
                                 llvm::SmallVectorImpl<Expr *> &SizeExprs);
   /// Parses a 'gang-arg-list', used for the 'gang' clause.
-  bool ParseOpenACCGangArgList(SourceLocation GangLoc);
-  /// Parses a 'gang-arg', used for the 'gang' clause.
-  bool ParseOpenACCGangArg(SourceLocation GangLoc);
+  bool ParseOpenACCGangArgList(SourceLocation GangLoc,
+                               llvm::SmallVectorImpl<OpenACCGangKind> &GKs,
+                               llvm::SmallVectorImpl<Expr *> &IntExprs);
+
+  using OpenACCGangArgRes = std::pair<OpenACCGangKind, ExprResult>;
+  /// Parses a 'gang-arg', used for the 'gang' clause. Returns a pair of the
+  /// ExprResult (which contains the validity of the expression), plus the gang
+  /// kind for the current argument.
+  OpenACCGangArgRes 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 97386d2378b758..59a9648d5f9380 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -38,9 +38,20 @@ class SemaOpenACC : public SemaBase {
   /// haven't had their 'parent' compute construct set yet. Entires will only be
   /// made to this list in the case where we know the loop isn't an orphan.
   llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
-  /// Whether we are inside of a compute construct, and should add loops to the
-  /// above collection.
-  bool InsideComputeConstruct = false;
+
+  struct ComputeConstructInfo {
+    /// Which type of compute construct we are inside of, which we can use to
+    /// determine whether we should add loops to the above collection.  We can
+    /// also use it to diagnose loop construct clauses.
+    OpenACCDirectiveKind Kind = OpenACCDirectiveKind::Invalid;
+    // If we have an active compute construct, stores the list of clauses we've
+    // prepared for it, so that we can diagnose limitations on child constructs.
+    ArrayRef<OpenACCClause *> Clauses;
+  } ActiveComputeConstructInfo;
+
+  bool isInComputeConstruct() const {
+    return ActiveComputeConstructInfo.Kind != OpenACCDirectiveKind::Invalid;
+  }
 
   /// Certain clauses care about the same things that aren't specific to the
   /// individual clause, but can be shared by a few, so store them here. All
@@ -99,6 +110,15 @@ class SemaOpenACC : public SemaBase {
   } TileInfo;
 
 public:
+  ComputeConstructInfo &getActiveComputeConstructInfo() {
+    return ActiveComputeConstructInfo;
+  }
+
+  /// If there is a current 'active' loop construct with a 'gang' clause on a
+  /// 'kernel' construct, this will have the source location for it. This
+  /// permits us to implement the restriction of no further 'gang' clauses.
+  SourceLocation LoopGangClauseOnKernelLoc;
+
   // Redeclaration of the version in OpenACCClause.h.
   using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
 
@@ -149,9 +169,14 @@ class SemaOpenACC : public SemaBase {
       Expr *LoopCount;
     };
 
+    struct GangDetails {
+      SmallVector<OpenACCGangKind> GangKinds;
+      SmallVector<Expr *> IntExprs;
+    };
+
     std::variant<std::monostate, DefaultDetails, ConditionDetails,
                  IntExprDetails, VarListDetails, WaitDetails, DeviceTypeDetails,
-                 ReductionDetails, CollapseDetails>
+                 ReductionDetails, CollapseDetails, GangDetails>
         Details = std::monostate{};
 
   public:
@@ -245,9 +270,18 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::Tile ||
+              ClauseKind == OpenACCClauseKind::Gang ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
 
+      if (ClauseKind == OpenACCClauseKind::Gang) {
+        // There might not be any gang int exprs, as this is an optional
+        // argument.
+        if (std::holds_alternative<std::monostate>(Details))
+          return {};
+        return std::get<GangDetails>(Details).IntExprs;
+      }
+
       return std::get<IntExprDetails>(Details).IntExprs;
     }
 
@@ -259,6 +293,16 @@ class SemaOpenACC : public SemaBase {
       return std::get<ReductionDetails>(Details).Op;
     }
 
+    ArrayRef<OpenACCGangKind> getGangKinds() const {
+      assert(ClauseKind == OpenACCClauseKind::Gang &&
+             "Parsed clause kind does not have gang kind");
+      // The args on gang are optional, so this might not actually hold
+      // anything.
+      if (std::holds_alternative<std::monostate>(Details))
+        return {};
+      return std::get<GangDetails>(Details).GangKinds;
+    }
+
     ArrayRef<Expr *> getVarList() {
       assert((ClauseKind == OpenACCClauseKind::Private ||
               ClauseKind == OpenACCClauseKind::NoCreate ||
@@ -371,6 +415,25 @@ class SemaOpenACC : public SemaBase {
       Details = IntExprDetails{std::move(IntExprs)};
     }
 
+    void setGangDetails(ArrayRef<OpenACCGangKind> GKs,
+                        ArrayRef<Expr *> IntExprs) {
+      assert(ClauseKind == OpenACCClauseKind::Gang &&
+             "Parsed Clause kind does not have gang details");
+      assert(GKs.size() == IntExprs.size() && "Mismatched kind/size?");
+
+      Details = GangDetails{{GKs.begin(), GKs.end()},
+                            {IntExprs.begin(), IntExprs.end()}};
+    }
+
+    void setGangDetails(llvm::SmallVector<OpenACCGangKind> &&GKs,
+                        llvm::SmallVector<Expr *> &&IntExprs) {
+      assert(ClauseKind == OpenACCClauseKind::Gang &&
+             "Parsed Clause kind does not have gang details");
+      assert(GKs.size() == IntExprs.size() && "Mismatched kind/size?");
+
+      Details = GangDetails{std::move(GKs), std::move(IntExprs)};
+    }
+
     void setVarListDetails(ArrayRef<Expr *> VarList, bool IsReadOnly,
                            bool IsZero) {
       assert((ClauseKind == OpenACCClauseKind::Private ||
@@ -545,10 +608,12 @@ class SemaOpenACC : public SemaBase {
                                    SourceLocation RBLoc);
   /// Checks the loop depth value for a collapse clause.
   ExprResult CheckCollapseLoopCount(Expr *LoopCount);
-  /// Checks a single size expr for a tile clause. 'gang' could possibly call
-  /// this, but has slightly stricter rules as to valid values.
+  /// Checks a single size expr for a tile clause.
   ExprResult CheckTileSizeExpr(Expr *SizeExpr);
 
+  // Check a single expression on a gang clause.
+  ExprResult CheckGangExpr(OpenACCGangKind GK, Expr *E);
+
   ExprResult BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc);
   ExprResult ActOnOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc);
 
@@ -595,8 +660,9 @@ class SemaOpenACC : public SemaBase {
   /// Loop needing its parent construct.
   class AssociatedStmtRAII {
     SemaOpenACC &SemaRef;
-    bool WasInsideComputeConstruct;
+    ComputeConstructInfo OldActiveComputeConstructInfo;
     OpenACCDirectiveKind DirKind;
+    SourceLocation OldLoopGangClauseOnKernelLoc;
     llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
     LoopInConstructRAII LoopRAII;
 

diff  --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 0b34ed6189593e..6fb8fe0b8cfeef 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -26,7 +26,7 @@ bool OpenACCClauseWithExprs::classof(const OpenACCClause *C) {
   return OpenACCWaitClause::classof(C) || OpenACCNumGangsClause::classof(C) ||
          OpenACCTileClause::classof(C) ||
          OpenACCClauseWithSingleIntExpr::classof(C) ||
-         OpenACCClauseWithVarList::classof(C);
+         OpenACCGangClause::classof(C) || OpenACCClauseWithVarList::classof(C);
 }
 bool OpenACCClauseWithVarList::classof(const OpenACCClause *C) {
   return OpenACCPrivateClause::classof(C) ||
@@ -125,6 +125,21 @@ OpenACCNumWorkersClause::OpenACCNumWorkersClause(SourceLocation BeginLoc,
          "Condition expression type not scalar/dependent");
 }
 
+OpenACCGangClause::OpenACCGangClause(SourceLocation BeginLoc,
+                                     SourceLocation LParenLoc,
+                                     ArrayRef<OpenACCGangKind> GangKinds,
+                                     ArrayRef<Expr *> IntExprs,
+                                     SourceLocation EndLoc)
+    : OpenACCClauseWithExprs(OpenACCClauseKind::Gang, BeginLoc, LParenLoc,
+                             EndLoc) {
+  assert(GangKinds.size() == IntExprs.size() && "Mismatch exprs/kind?");
+  std::uninitialized_copy(IntExprs.begin(), IntExprs.end(),
+                          getTrailingObjects<Expr *>());
+  setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), IntExprs.size()));
+  std::uninitialized_copy(GangKinds.begin(), GangKinds.end(),
+                          getTrailingObjects<OpenACCGangKind>());
+}
+
 OpenACCNumWorkersClause *
 OpenACCNumWorkersClause::Create(const ASTContext &C, SourceLocation BeginLoc,
                                 SourceLocation LParenLoc, Expr *IntExpr,
@@ -376,11 +391,16 @@ OpenACCSeqClause *OpenACCSeqClause::Create(const ASTContext &C,
   return new (Mem) OpenACCSeqClause(BeginLoc, EndLoc);
 }
 
-OpenACCGangClause *OpenACCGangClause::Create(const ASTContext &C,
-                                             SourceLocation BeginLoc,
-                                             SourceLocation EndLoc) {
-  void *Mem = C.Allocate(sizeof(OpenACCGangClause));
-  return new (Mem) OpenACCGangClause(BeginLoc, EndLoc);
+OpenACCGangClause *
+OpenACCGangClause::Create(const ASTContext &C, SourceLocation BeginLoc,
+                          SourceLocation LParenLoc,
+                          ArrayRef<OpenACCGangKind> GangKinds,
+                          ArrayRef<Expr *> IntExprs, SourceLocation EndLoc) {
+  void *Mem =
+      C.Allocate(OpenACCGangClause::totalSizeToAlloc<Expr *, OpenACCGangKind>(
+          IntExprs.size(), GangKinds.size()));
+  return new (Mem)
+      OpenACCGangClause(BeginLoc, LParenLoc, GangKinds, IntExprs, EndLoc);
 }
 
 OpenACCWorkerClause *OpenACCWorkerClause::Create(const ASTContext &C,
@@ -600,3 +620,21 @@ void OpenACCClausePrinter::VisitCollapseClause(const OpenACCCollapseClause &C) {
   printExpr(C.getLoopCount());
   OS << ")";
 }
+
+void OpenACCClausePrinter::VisitGangClause(const OpenACCGangClause &C) {
+  OS << "gang";
+
+  if (C.getNumExprs() > 0) {
+    OS << "(";
+    bool first = true;
+    for (unsigned I = 0; I < C.getNumExprs(); ++I) {
+      if (!first)
+        OS << ", ";
+      first = false;
+
+      OS << C.getExpr(I).first << ": ";
+      printExpr(C.getExpr(I).second);
+    }
+    OS << ")";
+  }
+}

diff  --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 4d177fd6c5968c..6161b1403ed35d 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2646,6 +2646,12 @@ void OpenACCClauseProfiler::VisitIndependentClause(
 
 void OpenACCClauseProfiler::VisitSeqClause(const OpenACCSeqClause &Clause) {}
 
+void OpenACCClauseProfiler::VisitGangClause(const OpenACCGangClause &Clause) {
+  for (unsigned I = 0; I < Clause.getNumExprs(); ++I) {
+    Profiler.VisitStmt(Clause.getExpr(I).second);
+  }
+}
+
 void OpenACCClauseProfiler::VisitReductionClause(
     const OpenACCReductionClause &Clause) {
   for (auto *E : Clause.getVarList())

diff  --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 15b23d60c3ffab..ac8c196777f9b8 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -425,6 +425,17 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
       // but print 'clause' here so it is clear what is happening from the dump.
       OS << " clause";
       break;
+    case OpenACCClauseKind::Gang: {
+      OS << " clause";
+      // print the list of all GangKinds, so that there is some sort of
+      // relationship to the expressions listed afterwards.
+      auto *GC = cast<OpenACCGangClause>(C);
+
+      for (unsigned I = 0; I < GC->getNumExprs(); ++I) {
+        OS << " " << GC->getExpr(I).first;
+      }
+      break;
+    }
     case OpenACCClauseKind::Collapse:
       OS << " clause";
       if (cast<OpenACCCollapseClause>(C)->hasForce())

diff  --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index b27e50b147f4a8..635039b724e6a0 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -797,23 +797,26 @@ bool Parser::ParseOpenACCSizeExprList(
 /// [num:]int-expr
 /// dim:int-expr
 /// static:size-expr
-bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
+Parser::OpenACCGangArgRes Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
 
   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Static, getCurToken()) &&
       NextToken().is(tok::colon)) {
     // 'static' just takes a size-expr, which is an int-expr or an asterisk.
     ConsumeToken();
     ConsumeToken();
-    return ParseOpenACCSizeExpr(OpenACCClauseKind::Gang).isInvalid();
+    ExprResult Res = ParseOpenACCSizeExpr(OpenACCClauseKind::Gang);
+    return {OpenACCGangKind::Static, Res};
   }
 
   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Dim, getCurToken()) &&
       NextToken().is(tok::colon)) {
     ConsumeToken();
     ConsumeToken();
-    return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
-                               OpenACCClauseKind::Gang, GangLoc)
-        .first.isInvalid();
+    // Parse this as a const-expression, and we'll check its integer-ness/value
+    // in CheckGangExpr.
+    ExprResult Res =
+        getActions().CorrectDelayedTyposInExpr(ParseConstantExpression());
+    return {OpenACCGangKind::Dim, Res};
   }
 
   if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Num, getCurToken()) &&
@@ -822,27 +825,40 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
     ConsumeToken();
     // Fallthrough to the 'int-expr' handling for when 'num' is omitted.
   }
+
   // This is just the 'num' case where 'num' is optional.
-  return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
-                             OpenACCClauseKind::Gang, GangLoc)
-      .first.isInvalid();
+  ExprResult Res = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
+                                       OpenACCClauseKind::Gang, GangLoc)
+                       .first;
+  return {OpenACCGangKind::Num, Res};
 }
 
-bool Parser::ParseOpenACCGangArgList(SourceLocation GangLoc) {
-  if (ParseOpenACCGangArg(GangLoc)) {
+bool Parser::ParseOpenACCGangArgList(
+    SourceLocation GangLoc, llvm::SmallVectorImpl<OpenACCGangKind> &GKs,
+    llvm::SmallVectorImpl<Expr *> &IntExprs) {
+
+  Parser::OpenACCGangArgRes Res = ParseOpenACCGangArg(GangLoc);
+  if (!Res.second.isUsable()) {
     SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
               Parser::StopBeforeMatch);
-    return false;
+    return true;
   }
 
+  GKs.push_back(Res.first);
+  IntExprs.push_back(Res.second.get());
+
   while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
     ExpectAndConsume(tok::comma);
 
-    if (ParseOpenACCGangArg(GangLoc)) {
+    Res = ParseOpenACCGangArg(GangLoc);
+    if (!Res.second.isUsable()) {
       SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
                 Parser::StopBeforeMatch);
-      return false;
+      return true;
     }
+
+    GKs.push_back(Res.first);
+    IntExprs.push_back(Res.second.get());
   }
   return false;
 }
@@ -1129,12 +1145,16 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
         }
         break;
       }
-      case OpenACCClauseKind::Gang:
-        if (ParseOpenACCGangArgList(ClauseLoc)) {
+      case OpenACCClauseKind::Gang: {
+        llvm::SmallVector<OpenACCGangKind> GKs;
+        llvm::SmallVector<Expr *> IntExprs;
+        if (ParseOpenACCGangArgList(ClauseLoc, GKs, IntExprs)) {
           Parens.skipToEnd();
           return OpenACCCanContinue();
         }
+        ParsedClause.setGangDetails(std::move(GKs), std::move(IntExprs));
         break;
+      }
       case OpenACCClauseKind::Wait: {
         OpenACCWaitParseInfo Info =
             ParseOpenACCWaitArgument(ClauseLoc,

diff  --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 66f8029a2754b9..30d73d621db69b 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -366,6 +366,19 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
     }
   }
 
+  case OpenACCClauseKind::Gang: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Loop:
+    case OpenACCDirectiveKind::ParallelLoop:
+    case OpenACCDirectiveKind::SerialLoop:
+    case OpenACCDirectiveKind::KernelsLoop:
+    case OpenACCDirectiveKind::Routine:
+      return true;
+    default:
+      return false;
+    }
+  }
+
   default:
     // Do nothing so we can go to the 'unimplemented' diagnostic instead.
     return true;
@@ -459,6 +472,23 @@ class SemaOpenACCClauseVisitor {
     return nullptr;
   }
 
+  // OpenACC 3.3 2.9:
+  // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause
+  // appears.
+  bool DiagIfSeqClause(SemaOpenACC::OpenACCParsedClause &Clause) {
+    const auto *Itr =
+        llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCSeqClause>);
+
+    if (Itr != ExistingClauses.end()) {
+      SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine)
+          << Clause.getClauseKind() << (*Itr)->getClauseKind();
+      SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+
+      return true;
+    }
+    return false;
+  }
+
 public:
   SemaOpenACCClauseVisitor(SemaOpenACC &S,
                            ArrayRef<const OpenACCClause *> ExistingClauses)
@@ -470,26 +500,14 @@ class SemaOpenACCClauseVisitor {
 
   OpenACCClause *Visit(SemaOpenACC::OpenACCParsedClause &Clause) {
     switch (Clause.getClauseKind()) {
-  case OpenACCClauseKind::Gang:
-  case OpenACCClauseKind::Worker:
-  case OpenACCClauseKind::Vector: {
-    // TODO OpenACC: These are only implemented enough for the 'seq' diagnostic,
-    // otherwise treats itself as unimplemented.  When we implement these, we
-    // can remove them from here.
-
-    // OpenACC 3.3 2.9:
-    // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause
-    // appears.
-    const auto *Itr =
-        llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCSeqClause>);
-
-    if (Itr != ExistingClauses.end()) {
-      SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine)
-          << Clause.getClauseKind() << (*Itr)->getClauseKind();
-      SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+    case OpenACCClauseKind::Worker:
+    case OpenACCClauseKind::Vector: {
+      // TODO OpenACC: These are only implemented enough for the 'seq'
+      // diagnostic, otherwise treats itself as unimplemented.  When we
+      // implement these, we can remove them from here.
+      DiagIfSeqClause(Clause);
+      return isNotImplemented();
     }
-    return isNotImplemented();
-  }
 
 #define VISIT_CLAUSE(CLAUSE_NAME)                                              \
   case OpenACCClauseKind::CLAUSE_NAME:                                         \
@@ -1006,6 +1024,84 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause(
                                           Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
+    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();
+
+  llvm::SmallVector<OpenACCGangKind> GangKinds;
+  llvm::SmallVector<Expr *> IntExprs;
+
+  // Store the existing locations, so we can do duplicate checking.  Index is
+  // the int-value of the OpenACCGangKind enum.
+  SourceLocation ExistingElemLoc[3];
+
+  for (unsigned I = 0; I < Clause.getIntExprs().size(); ++I) {
+    OpenACCGangKind GK = Clause.getGangKinds()[I];
+    ExprResult ER = SemaRef.CheckGangExpr(GK, Clause.getIntExprs()[I]);
+
+    if (!ER.isUsable())
+      continue;
+
+    // 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
+    // keyword or with num keyword is only allowed when num_gangs does not
+    // appear on the kernels construct.
+    if (SemaRef.getActiveComputeConstructInfo().Kind ==
+            OpenACCDirectiveKind::Kernels &&
+        GK == OpenACCGangKind::Num) {
+      const auto *Itr =
+          llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
+                        llvm::IsaPred<OpenACCNumGangsClause>);
+
+      if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
+        SemaRef.Diag(ER.get()->getBeginLoc(),
+                     diag::err_acc_gang_num_gangs_conflict);
+        SemaRef.Diag((*Itr)->getBeginLoc(),
+                     diag::note_acc_previous_clause_here);
+        continue;
+      }
+    }
+
+    // OpenACC 3.3 2.9: 'gang-arg-list' may have at most one num, one dim, and
+    // one static argument.
+    if (ExistingElemLoc[static_cast<unsigned>(GK)].isValid()) {
+      SemaRef.Diag(ER.get()->getBeginLoc(), diag::err_acc_gang_multiple_elt)
+          << static_cast<unsigned>(GK);
+      SemaRef.Diag(ExistingElemLoc[static_cast<unsigned>(GK)],
+                   diag::note_acc_previous_expr_here);
+      continue;
+    }
+
+    ExistingElemLoc[static_cast<unsigned>(GK)] = ER.get()->getBeginLoc();
+    GangKinds.push_back(GK);
+    IntExprs.push_back(ER.get());
+  }
+
+  // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+  // construct, the gang clause behaves as follows. ... The region of a loop
+  // with a gang clause may not contain another loop with a gang clause unless
+  // within a nested compute region.
+  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(SemaRef.LoopGangClauseOnKernelLoc,
+                 diag::note_acc_previous_clause_here);
+    return nullptr;
+  }
+
+  return OpenACCGangClause::Create(Ctx, Clause.getBeginLoc(),
+                                   Clause.getLParenLoc(), GangKinds, IntExprs,
+                                   Clause.getEndLoc());
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   // Restrictions only properly implemented on 'loop' constructs, and it is
@@ -1118,17 +1214,44 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     SemaOpenACC &S, OpenACCDirectiveKind DK,
     ArrayRef<const OpenACCClause *> UnInstClauses,
     ArrayRef<OpenACCClause *> Clauses)
-    : SemaRef(S), WasInsideComputeConstruct(S.InsideComputeConstruct),
-      DirKind(DK), LoopRAII(SemaRef, /*PreserveDepth=*/false) {
+    : SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo),
+      DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc),
+      LoopRAII(SemaRef, /*PreserveDepth=*/false) {
   // Compute constructs end up taking their 'loop'.
   if (DirKind == OpenACCDirectiveKind::Parallel ||
       DirKind == OpenACCDirectiveKind::Serial ||
       DirKind == OpenACCDirectiveKind::Kernels) {
-    SemaRef.InsideComputeConstruct = true;
+    SemaRef.ActiveComputeConstructInfo.Kind = DirKind;
+    SemaRef.ActiveComputeConstructInfo.Clauses = Clauses;
     SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
+
+    // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+    // construct, the gang clause behaves as follows. ... The region of a loop
+    // with a gang clause may not contain another loop with a gang clause unless
+    // within a nested compute region.
+    //
+    // Implement the 'unless within a nested compute region' part.
+    SemaRef.LoopGangClauseOnKernelLoc = {};
   } else if (DirKind == OpenACCDirectiveKind::Loop) {
     SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
     SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
+
+    // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+    // construct, the gang clause behaves as follows. ... The region of a loop
+    // with a gang clause may not contain another loop with a gang clause unless
+    // within a nested compute region.
+    //
+    // We don't bother doing this when this is a template instantiation, as
+    // there is no reason to do these checks: the existance of a
+    // gang/kernels/etc cannot be dependent.
+    if (SemaRef.getActiveComputeConstructInfo().Kind ==
+            OpenACCDirectiveKind::Kernels &&
+        UnInstClauses.empty()) {
+      // This handles the 'outer loop' part of this.
+      auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCGangClause>);
+      if (Itr != Clauses.end())
+        SemaRef.LoopGangClauseOnKernelLoc = (*Itr)->getBeginLoc();
+    }
   }
 }
 
@@ -1199,7 +1322,9 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
 }
 
 SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
-  SemaRef.InsideComputeConstruct = WasInsideComputeConstruct;
+  SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
+  SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc;
+
   if (DirKind == OpenACCDirectiveKind::Parallel ||
       DirKind == OpenACCDirectiveKind::Serial ||
       DirKind == OpenACCDirectiveKind::Kernels) {
@@ -1761,6 +1886,109 @@ ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr *LoopCount) {
       ConstantExpr::Create(getASTContext(), LoopCount, APValue{*ICE})};
 }
 
+namespace {
+ExprResult CheckGangStaticExpr(SemaOpenACC &S, Expr *E) {
+  if (isa<OpenACCAsteriskSizeExpr>(E))
+    return E;
+  return S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang,
+                        E->getBeginLoc(), E);
+}
+} // namespace
+
+ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK, Expr *E) {
+  // Gang Expr legality depends on the associated compute construct.
+  switch (ActiveComputeConstructInfo.Kind) {
+  case OpenACCDirectiveKind::Invalid:
+  case OpenACCDirectiveKind::Parallel: {
+    switch (GK) {
+      // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel
+      // construct, or an orphaned loop construct, the gang clause behaves as
+      // follows. ... The dim argument must be a constant positive integer value
+      // 1, 2, or 3.
+    case OpenACCGangKind::Dim: {
+      if (!E)
+        return ExprError();
+      ExprResult Res =
+          ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang,
+                       E->getBeginLoc(), E);
+
+      if (!Res.isUsable())
+        return Res;
+
+      if (Res.get()->isInstantiationDependent())
+        return Res;
+
+      std::optional<llvm::APSInt> ICE =
+          Res.get()->getIntegerConstantExpr(getASTContext());
+
+      if (!ICE || *ICE <= 0 || ICE > 3) {
+        Diag(Res.get()->getBeginLoc(), diag::err_acc_gang_dim_value)
+            << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue();
+        return ExprError();
+      }
+
+      return ExprResult{
+          ConstantExpr::Create(getASTContext(), Res.get(), APValue{*ICE})};
+    }
+      // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel
+      // 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
+          << (/*orphan/parallel=*/ActiveComputeConstructInfo.Kind ==
+                      OpenACCDirectiveKind::Parallel
+                  ? 1
+                  : 0);
+      return ExprError();
+    case OpenACCGangKind::Static:
+      return CheckGangStaticExpr(*this, E);
+    }
+  } break;
+  case OpenACCDirectiveKind::Kernels: {
+    switch (GK) {
+    // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+    // 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;
+      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
+    // keyword or with num keyword is only allowed when num_gangs does not
+    // appear on the kernels construct. ... The region of a loop with the gang
+    // clause may not contain another loop with a gang clause unless within a
+    // nested compute region.
+    case OpenACCGangKind::Num:
+      // This isn't allowed if there is a 'num_gangs' on the kernel construct,
+      // and makes loop-with-gang-clause ill-formed inside of this 'loop', but
+      // nothing can be enforced here.
+      return ExprResult{E};
+    case OpenACCGangKind::Static:
+      return CheckGangStaticExpr(*this, E);
+    }
+  } break;
+  case OpenACCDirectiveKind::Serial: {
+    switch (GK) {
+    // 'dim' and 'num' don't really make sense on serial, and GCC rejects them
+    // too, so we disallow them too.
+    case OpenACCGangKind::Dim:
+    case OpenACCGangKind::Num:
+      Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid)
+          << GK << /*Kernels=*/3;
+      return ExprError();
+    case OpenACCGangKind::Static:
+      return CheckGangStaticExpr(*this, E);
+    }
+  }
+  default:
+    llvm_unreachable("Non compute construct in active compute construct?");
+  }
+
+  llvm_unreachable("Compute construct directive not handled?");
+}
+
 ExprResult SemaOpenACC::CheckTileSizeExpr(Expr *SizeExpr) {
   if (!SizeExpr)
     return ExprError();
@@ -2031,7 +2259,7 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(OpenACCDirectiveKind K,
     // If we are in the scope of a compute construct, add this to the list of
     // loop constructs that need assigning to the next closing compute
     // construct.
-    if (InsideComputeConstruct)
+    if (isInComputeConstruct())
       ParentlessLoopConstructs.push_back(LoopConstruct);
 
     return LoopConstruct;

diff  --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 5753c9eccf6c92..cde40773336866 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11940,6 +11940,29 @@ void OpenACCClauseTransform<Derived>::VisitTileClause(
       ParsedClause.getLParenLoc(), ParsedClause.getIntExprs(),
       ParsedClause.getEndLoc());
 }
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitGangClause(
+    const OpenACCGangClause &C) {
+  llvm::SmallVector<OpenACCGangKind> TransformedGangKinds;
+  llvm::SmallVector<Expr *> TransformedIntExprs;
+
+  for (unsigned I = 0; I < C.getNumExprs(); ++I) {
+    ExprResult ER = Self.TransformExpr(const_cast<Expr *>(C.getExpr(I).second));
+    if (!ER.isUsable())
+      continue;
+
+    ER = Self.getSema().OpenACC().CheckGangExpr(C.getExpr(I).first, ER.get());
+    if (!ER.isUsable())
+      continue;
+    TransformedGangKinds.push_back(C.getExpr(I).first);
+    TransformedIntExprs.push_back(ER.get());
+  }
+
+  NewClause = OpenACCGangClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), TransformedGangKinds, TransformedIntExprs,
+      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 e638129897692f..0339419da43cab 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12326,6 +12326,18 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCTileClause::Create(getContext(), BeginLoc, LParenLoc,
                                      SizeExprs, EndLoc);
   }
+  case OpenACCClauseKind::Gang: {
+    SourceLocation LParenLoc = readSourceLocation();
+    unsigned NumExprs = readInt();
+    llvm::SmallVector<OpenACCGangKind> GangKinds;
+    llvm::SmallVector<Expr *> Exprs;
+    for (unsigned I = 0; I < NumExprs; ++I) {
+      GangKinds.push_back(readEnum<OpenACCGangKind>());
+      Exprs.push_back(readSubExpr());
+    }
+    return OpenACCGangClause::Create(getContext(), BeginLoc, LParenLoc,
+                                     GangKinds, Exprs, EndLoc);
+  }
 
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
@@ -12342,7 +12354,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   case OpenACCClauseKind::Bind:
   case OpenACCClauseKind::DeviceNum:
   case OpenACCClauseKind::DefaultAsync:
-  case OpenACCClauseKind::Gang:
   case OpenACCClauseKind::Invalid:
     llvm_unreachable("Clause serialization not yet implemented");
   }

diff  --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 4976327fc654ee..583d9a4bccb800 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8182,6 +8182,16 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
       AddStmt(E);
     return;
   }
+  case OpenACCClauseKind::Gang: {
+    const auto *GC = cast<OpenACCGangClause>(C);
+    writeSourceLocation(GC->getLParenLoc());
+    writeUInt32(GC->getNumExprs());
+    for (unsigned I = 0; I < GC->getNumExprs(); ++I) {
+      writeEnum(GC->getExpr(I).first);
+      AddStmt(const_cast<Expr *>(GC->getExpr(I).second));
+    }
+    return;
+  }
 
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
@@ -8198,7 +8208,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::Bind:
   case OpenACCClauseKind::DeviceNum:
   case OpenACCClauseKind::DefaultAsync:
-  case OpenACCClauseKind::Gang:
   case OpenACCClauseKind::Invalid:
     llvm_unreachable("Clause serialization not yet implemented");
   }

diff  --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp
index aee4591cab428f..baa4b173f88edc 100644
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -95,4 +95,86 @@ void foo() {
     for(;;)
       for(;;)
         for(;;);
+
+// CHECK: #pragma acc loop gang(dim: 2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop gang(dim:2)
+  for(;;);
+
+// CHECK: #pragma acc loop gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop gang(static:i)
+  for(;;);
+
+// CHECK: #pragma acc loop gang(static: i) gang(dim: 2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop gang(static:i) gang(dim:2)
+  for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop gang(dim: 2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop gang(dim:2)
+  for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop gang(static:i)
+  for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop gang(static: i) gang(dim: 2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop gang(static:i) gang(dim:2)
+  for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop gang(num: i) gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop gang(i) gang(static:i)
+  for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop gang(num: i) gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop gang(num:i) gang(static:i)
+  for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop gang(static:i)
+  for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop gang(static: *)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop gang(static:*)
+  for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop gang
+  for(;;);
 }

diff  --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 6c382379a8a7ea..899fbd78b87298 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -1202,7 +1202,6 @@ void Tile() {
 }
 
 void Gang() {
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang
   for(;;){}
   // expected-error at +3{{expected expression}}
@@ -1210,68 +1209,58 @@ void Gang() {
   // expected-note at +1{{to match this '('}}
 #pragma acc loop gang(
   for(;;){}
-  // expected-error at +2{{expected expression}}
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +1{{expected expression}}
 #pragma acc loop gang()
   for(;;){}
 
-  // expected-error at +2{{expected expression}}
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +1{{expected expression}}
 #pragma acc loop gang(5, *)
   for(;;){}
 
-  // expected-error at +2{{expected expression}}
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +1{{expected expression}}
 #pragma acc loop gang(*)
   for(;;){}
 
-  // expected-error at +2{{expected expression}}
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +1{{expected expression}}
 #pragma acc loop gang(5, num:*)
   for(;;){}
 
-  // expected-error at +2{{expected expression}}
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +1{{expected expression}}
 #pragma acc loop gang(num:5, *)
   for(;;){}
 
-  // expected-error at +2{{expected expression}}
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +1{{expected expression}}
 #pragma acc loop gang(num:5, num:*)
   for(;;){}
 
-  // expected-error at +2{{expected expression}}
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +1{{expected expression}}
 #pragma acc loop gang(num:*)
   for(;;){}
 
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-#pragma acc loop gang(dim:5)
+#pragma acc loop gang(dim:2)
   for(;;){}
 
-  // expected-error at +2{{expected expression}}
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +1{{expected expression}}
 #pragma acc loop gang(dim:5, dim:*)
   for(;;){}
 
-  // expected-error at +2{{expected expression}}
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +1{{expected expression}}
 #pragma acc loop gang(dim:*)
   for(;;){}
 
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(static:*)
   for(;;){}
 
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +2{{OpenACC 'gang' clause may have at most one 'static' argument}}
+  // expected-note at +1{{previous expression is here}}
 #pragma acc loop gang(static:*, static:5)
   for(;;){}
 
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+#pragma acc kernels
 #pragma acc loop gang(static:*, 5)
   for(;;){}
 
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+#pragma acc kernels
 #pragma acc loop gang(static:45, 5)
   for(;;){}
 
@@ -1330,11 +1319,16 @@ void Gang() {
 #pragma acc loop gang(dim:45
   for(;;){}
 
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-#pragma acc loop gang(static:*, dim:returns_int(), 5)
+#pragma acc kernels
+#pragma acc loop gang(static:*, 5)
+  for(;;){}
+
+  // expected-error at +1{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc loop gang(static:*, dim:returns_int())
   for(;;){}
 
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+  // expected-error at +2 2{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+  // expected-error at +1{{argument to 'gang' clause dimension must be a constant expression}}
 #pragma acc loop gang(num: 32, static:*, dim:returns_int(), 5)
   for(;;){}
 

diff  --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
index d08497a7782edb..89000517c43fb5 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -214,8 +214,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'tile' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) tile(Var, 1)
   while(1);
-  // expected-error at +2{{OpenACC clause 'gang' may not follow a 'dtype' clause in a compute construct}}
-  // expected-note at +1{{previous clause is here}}
+  // expected-error at +1{{OpenACC 'gang' clause is not valid on 'kernels' directive}}
 #pragma acc kernels dtype(*) gang
   while(1);
 #pragma acc kernels device_type(*) wait

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 3da7f0e9836be8..6c2c79b02a4131 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -168,7 +168,6 @@ void uses() {
 #pragma acc loop auto tile(1+2, 1)
   for(;;)
     for(;;);
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc loop auto gang
   for(;;);
   // expected-error at +1{{OpenACC 'wait' clause is not valid on 'loop' directive}}
@@ -306,7 +305,6 @@ void uses() {
 #pragma acc loop tile(1+2, 1) auto
   for(;;)
     for(;;);
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc loop gang auto
   for(;;);
   // expected-error at +1{{OpenACC 'wait' clause is not valid on 'loop' directive}}
@@ -445,7 +443,6 @@ void uses() {
 #pragma acc loop independent tile(1+2, 1)
   for(;;)
     for(;;);
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc loop independent gang
   for(;;);
   // expected-error at +1{{OpenACC 'wait' clause is not valid on 'loop' directive}}
@@ -583,7 +580,6 @@ void uses() {
 #pragma acc loop tile(1+2, 1) independent
   for(;;)
     for(;;);
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc loop gang independent
   for(;;);
   // expected-error at +1{{OpenACC 'wait' clause is not valid on 'loop' directive}}
@@ -591,9 +587,8 @@ void uses() {
   for(;;);
 
   // 'seq' cannot be combined with 'gang', 'worker' or 'vector'
-  // expected-error at +3{{OpenACC clause 'gang' 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 'gang' not yet implemented}}
+  // expected-error at +2{{OpenACC clause 'gang' 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 gang
   for(;;);
   // expected-error at +3{{OpenACC clause 'worker' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
@@ -735,10 +730,8 @@ void uses() {
 #pragma acc loop seq wait
   for(;;);
 
-  // TODO OpenACC: when 'gang' is implemented and makes it to the AST, this should diagnose because of a conflict with 'seq'.
-  // TODOexpected-error at +3{{OpenACC clause 'gang' 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 'gang' not yet implemented}}
+  // expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'loop' construct}}
+  // 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'.

diff  --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index 3d77c031f42630..cedef3ca858f5e 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -193,7 +193,6 @@ void uses() {
   for(;;)
     for(;;);
 
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop dtype(*) gang
   for(;;);
   // expected-error at +1{{OpenACC 'wait' clause is not valid on 'loop' directive}}

diff  --git a/clang/test/SemaOpenACC/loop-construct-gang-ast.cpp b/clang/test/SemaOpenACC/loop-construct-gang-ast.cpp
new file mode 100644
index 00000000000000..e797d842e240dc
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-gang-ast.cpp
@@ -0,0 +1,330 @@
+// 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
+
+void NormalUses() {
+  // CHECK: FunctionDecl{{.*}}NormalUses
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: gang clause dim
+  // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop gang(dim:1)
+  for(;;);
+
+  int Val;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} used Val 'int'
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop gang(static:Val)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause num
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' '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 gang(num:1) gang(static:Val)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause dim static
+  // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop gang(dim:1, static:Val)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop gang(static:Val)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: OpenACCAsteriskSizeExpr
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop gang(static:*)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang 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 gang
+  for(;;);
+}
+
+template<typename T, unsigned One>
+void TemplateUses(T Val) {
+  // CHECK: FunctionTemplateDecl{{.*}}TemplateUses
+  // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 One
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplateUses 'void (T)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} referenced Val 'T'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: gang clause dim
+  // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}} 'One' 'unsigned int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop gang(dim:One)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop gang(static:Val)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: OpenACCAsteriskSizeExpr
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop gang(static:*)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause dim
+  // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}} 'One' 'unsigned int'
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop gang(dim:One) gang(static:Val)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause dim static
+  // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}} 'One' 'unsigned int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop gang(dim:One, static:Val)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop gang(static:Val)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang 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 gang
+  for(;;);
+
+  // Instantiation:
+  // CHECK-NEXT: FunctionDecl{{.*}} used TemplateUses 'void (int)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'int'
+  // CHECK-NEXT: BuiltinType{{.*}} 'int'
+  // CHECK-NEXT: TemplateArgument integral '1U'
+  // CHECK-NEXT: ParmVarDecl{{.*}} used Val 'int'
+  // CHECK-NEXT: CompoundStmt
+  //
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: gang clause dim
+  // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 One
+  // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: OpenACCAsteriskSizeExpr
+  // 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: gang clause dim
+  // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 One
+  // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' '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]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause dim static
+  // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 One
+  // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' '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]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' '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]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: gang clause
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: NullStmt
+}
+
+void inst() {
+  TemplateUses<int, 1>(5);
+}
+
+#endif // PCH_HELPER

diff  --git a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
new file mode 100644
index 00000000000000..ab6439ae576193
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
@@ -0,0 +1,335 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S{};
+struct Converts{
+  operator int();
+};
+
+template<typename T, unsigned Zero, unsigned Two, unsigned Four>
+void ParallelOrOrphanTempl() {
+  T i;
+  // expected-error at +1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop gang(i)
+  for(;;);
+  // expected-error at +1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop gang(num:i)
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop gang(i)
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop gang(num:i)
+  for(;;);
+
+  // expected-error at +1{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc loop gang(dim:i)
+  for(;;);
+
+  // expected-error at +2{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc parallel
+#pragma acc loop gang(dim:i)
+  for(;;);
+
+  // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc loop gang(dim:Zero)
+  for(;;);
+
+  // expected-error at +2{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc parallel
+#pragma acc loop gang(dim:Zero)
+  for(;;);
+
+  // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc loop gang(dim:Four)
+  for(;;);
+
+  // expected-error at +2{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc parallel
+#pragma acc loop gang(dim:Four)
+  for(;;);
+
+#pragma acc loop gang(static:i) gang(dim:Two)
+  for(;;);
+
+#pragma acc parallel
+#pragma acc loop gang(dim:Two) gang(static:*)
+  for(;;);
+
+#pragma acc parallel
+#pragma acc loop gang(dim:Two, static:i)
+  for(;;);
+
+  // expected-error at +4{{OpenACC 'gang' clause may have at most one 'static' argument}}
+  // expected-note at +3{{previous expression is here}}
+  // expected-error at +2{{OpenACC 'gang' clause may have at most one 'dim' argument}}
+  // expected-note at +1{{previous expression is here}}
+#pragma acc loop gang(static:i, static:i, dim:Two, dim:1)
+  for(;;);
+}
+
+void ParallelOrOrphan() {
+  ParallelOrOrphanTempl<int, 0, 2, 4>(); // expected-note{{in instantiation of function template}}
+
+  int i;
+  // expected-error at +1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop gang(i)
+  for(;;);
+  // expected-error at +1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop gang(num:i)
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop gang(i)
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop gang(num:i)
+  for(;;);
+
+  // expected-error at +1{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc loop gang(dim:i)
+  for(;;);
+
+  // expected-error at +2{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc parallel
+#pragma acc loop gang(dim:i)
+  for(;;);
+
+  // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc loop gang(dim:0)
+  for(;;);
+
+  // expected-error at +2{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc parallel
+#pragma acc loop gang(dim:0)
+  for(;;);
+
+  // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc loop gang(dim:4)
+  for(;;);
+
+  // expected-error at +2{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc parallel
+#pragma acc loop gang(dim:4)
+  for(;;);
+
+#pragma acc loop gang(static:i) gang(dim:2)
+  for(;;);
+
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:i)
+  for(;;);
+
+  S s;
+  // expected-error at +2{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}}
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:s)
+  for(;;);
+
+  Converts C;
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:C)
+  for(;;);
+}
+
+template<typename SomeS, typename SomeC, typename Int>
+void StaticIsIntegralTempl() {
+  SomeS s;
+  // expected-error at +2{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}}
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:s)
+  for(;;);
+
+  SomeC C;
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:C)
+  for(;;);
+  Int i;
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:i)
+  for(;;);
+
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:*)
+  for(;;);
+}
+
+void StaticIsIntegral() {
+  StaticIsIntegralTempl<S, Converts, int>();// expected-note{{in instantiation of function template}}
+
+  S s;
+  // expected-error at +2{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}}
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:s)
+  for(;;);
+
+  Converts C;
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:C)
+  for(;;);
+}
+
+template<unsigned I>
+void SerialTempl() {
+  // expected-error at +2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(I)
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(num:I)
+  for(;;);
+
+  // expected-error at +2{{'dim' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(dim:I)
+  for(;;);
+
+#pragma acc serial
+#pragma acc loop gang(static:I)
+  for(;;);
+}
+
+void Serial() {
+  SerialTempl<2>();
+
+  // expected-error at +2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(1)
+  for(;;);
+
+  // expected-error at +2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(num:1)
+  for(;;);
+
+  // expected-error at +2{{'dim' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(dim:1)
+  for(;;);
+
+#pragma acc serial
+#pragma acc loop gang(static:1)
+  for(;;);
+
+  int i;
+
+#pragma acc serial
+#pragma acc loop gang(static:i)
+  for(;;);
+}
+
+template<typename T>
+void KernelsTempl() {
+  T t;
+  // expected-error at +2{{'dim' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'kernels' compute construct}}
+#pragma acc kernels
+#pragma acc loop gang(dim:t)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop gang(static:t)
+  for(;;);
+
+  // expected-error at +3{{'num' argument to 'gang' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_gangs' clause}}
+  // expected-note at +1{{previous clause is here}}
+#pragma acc kernels num_gangs(t)
+#pragma acc loop gang(t)
+  for(;;);
+
+  // expected-error at +3{{'num' argument to 'gang' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_gangs' clause}}
+  // expected-note at +1{{previous clause is here}}
+#pragma acc kernels num_gangs(t)
+#pragma acc loop gang(num:t)
+  for(;;);
+}
+
+void Kernels() {
+  KernelsTempl<unsigned>();
+
+  // expected-error at +2{{'dim' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'kernels' compute construct}}
+#pragma acc kernels
+#pragma acc loop gang(dim:1)
+  for(;;);
+  unsigned t;
+#pragma acc kernels
+#pragma acc loop gang(static:t)
+  for(;;);
+
+  // expected-error at +3{{'num' argument to 'gang' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_gangs' clause}}
+  // expected-note at +1{{previous clause is here}}
+#pragma acc kernels num_gangs(1)
+#pragma acc loop gang(1)
+  for(;;);
+
+  // expected-error at +3{{'num' argument to 'gang' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_gangs' clause}}
+  // expected-note at +1{{previous clause is here}}
+#pragma acc kernels num_gangs(1)
+#pragma acc loop gang(num:1)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop gang(num:1)
+  for(;;) {
+    // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels' compute construct}}
+    // expected-note at -3{{previous clause is here}}
+#pragma acc loop gang(static:1)
+    for(;;);
+  }
+
+#pragma acc kernels
+#pragma acc loop gang(num:1)
+  for(;;) {
+    // allowed, intervening compute construct
+#pragma acc serial
+#pragma acc loop gang(static:1)
+    for(;;);
+  }
+
+#pragma acc kernels
+#pragma acc loop gang(num:1)
+  for(;;);
+
+  // OK, on a 
diff erent 'loop', not in the assoc statement.
+#pragma acc loop gang(static:1)
+  for(;;);
+
+  // expected-error at +3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}}
+  // expected-note at +2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(5, num:1)
+  for(;;);
+
+  // expected-error at +3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}}
+  // expected-note at +2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(num:5, 1)
+  for(;;);
+
+  // expected-error at +3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}}
+  // expected-note at +2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(num:5, num:1)
+  for(;;);
+}
+
+void MaxOneEntry() {
+  // expected-error at +3{{OpenACC 'gang' clause may have at most one 'static' argument}}
+  // expected-note at +2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(static: 1, static:1)
+    for(;;);
+
+#pragma acc kernels
+#pragma acc loop gang gang(static:1)
+    for(;;);
+}
+
+

diff  --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index c282a9071391e7..2ffe47fbd74476 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2907,6 +2907,11 @@ void OpenACCClauseEnqueue::VisitSeqClause(const OpenACCSeqClause &C) {}
 void OpenACCClauseEnqueue::VisitCollapseClause(const OpenACCCollapseClause &C) {
   Visitor.AddStmt(C.getLoopCount());
 }
+void OpenACCClauseEnqueue::VisitGangClause(const OpenACCGangClause &C) {
+  for (unsigned I = 0; I < C.getNumExprs(); ++I) {
+    Visitor.AddStmt(C.getExpr(I).second);
+  }
+}
 } // namespace
 
 void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {


        


More information about the cfe-commits mailing list