[clang] 3a4b9f3 - [OpenACC] Implement 'gang' clause for Combined Constructs

via cfe-commits cfe-commits at lists.llvm.org
Thu Dec 5 06:35:42 PST 2024


Author: erichkeane
Date: 2024-12-05T06:35:36-08:00
New Revision: 3a4b9f38915625c68c78b62de48a3de8b97c5043

URL: https://github.com/llvm/llvm-project/commit/3a4b9f38915625c68c78b62de48a3de8b97c5043
DIFF: https://github.com/llvm/llvm-project/commit/3a4b9f38915625c68c78b62de48a3de8b97c5043.diff

LOG: [OpenACC] Implement 'gang' clause for Combined Constructs

This one is a bit complicated, as it has some interesting interactions,
as 'gang' Sema is required to look at its containing compute construct.
Except in the case of a combined construct, they are the same. This
resulted in a large refactor of the checking code for CheckGangExpr,
plus some additional work on the diagnostics for its interaction with
'num_gangs' and 'vector'/'worker'.

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

Modified: 
    clang/include/clang/AST/OpenACCClause.h
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/include/clang/Sema/SemaOpenACC.h
    clang/lib/Sema/SemaOpenACC.cpp
    clang/lib/Sema/TreeTransform.h
    clang/test/AST/ast-print-openacc-combined-construct.cpp
    clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
    clang/test/SemaOpenACC/combined-construct-device_type-clause.c
    clang/test/SemaOpenACC/loop-construct-gang-clause.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 5ad4c336b6c531..2588c3f645c02b 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -483,6 +483,14 @@ class OpenACCGangClause final
     return {getGangKind(I), getExprs()[I]};
   }
 
+  bool hasExprOfKind(OpenACCGangKind GK) const {
+    for (unsigned I = 0; I < getNumExprs(); ++I) {
+      if (getGangKind(I) == GK)
+        return true;
+    }
+    return false;
+  }
+
   static OpenACCGangClause *
   Create(const ASTContext &Ctx, SourceLocation BeginLoc,
          SourceLocation LParenLoc, ArrayRef<OpenACCGangKind> GangKinds,

diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 2137cb713164ad..447358f0a5f382 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12755,21 +12755,22 @@ 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_int_arg_invalid
-    : Error<"'%1' argument on '%0' clause is not permitted on a%select{n "
-            "orphaned|||}2 'loop' construct %select{|associated with a "
-            "'parallel' compute construct|associated with a 'kernels' compute "
-            "construct|associated with a 'serial' compute construct}2">;
+    : Error<"'%0' argument on '%1' clause is not permitted on a%select{|n "
+            "orphaned}2 '%3' construct%select{| associated with a '%5' compute "
+            "construct}4">;
 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_num_arg_conflict
-    : Error<"'num' argument to '%0' clause not allowed on a 'loop' construct "
-            "associated with a 'kernels' construct that has a "
-            "'%select{num_gangs|num_workers|vector_length}1' "
-            "clause">;
+    : Error<"'num' argument to '%0' clause not allowed on a '%1' "
+            "construct%select{| associated with a '%3' construct}2 that has a "
+            "'%4' clause">;
+def err_acc_num_arg_conflict_reverse
+    : Error<"'num_gangs' clause not allowed on a 'kernels loop' construct that "
+            "has a 'gang' clause with a 'num' argument">;
 def err_acc_clause_in_clause_region
     : Error<"loop with a '%0' clause may not exist in the region of a '%1' "
-            "clause%select{| on a 'kernels' compute construct}2">;
+            "clause%select{| on a '%3' construct}2">;
 def err_acc_gang_reduction_conflict
     : Error<"%select{OpenACC 'gang' clause with a 'dim' value greater than "
             "1|OpenACC 'reduction' clause}0 cannot "

diff  --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index d720cf3c74d87e..a4132534686a81 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -164,9 +164,14 @@ class SemaOpenACC : public SemaBase {
   }
 
   /// 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;
+  /// 'kernel' construct, this will have the source location for it, and the
+  /// 'kernel kind'. This permits us to implement the restriction of no further
+  /// 'gang' clauses.
+  struct LoopGangOnKernelTy {
+    SourceLocation Loc;
+    OpenACCDirectiveKind DirKind = OpenACCDirectiveKind::Invalid;
+  } LoopGangClauseOnKernel;
+
   /// If there is a current 'active' loop construct with a 'worker' clause on it
   /// (on any sort of construct), this has the source location for it.  This
   /// permits us to implement the restriction of no further 'gang' or 'worker'
@@ -705,7 +710,9 @@ class SemaOpenACC : public SemaBase {
   ExprResult CheckTileSizeExpr(Expr *SizeExpr);
 
   // Check a single expression on a gang clause.
-  ExprResult CheckGangExpr(OpenACCGangKind GK, Expr *E);
+  ExprResult CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses,
+                           OpenACCDirectiveKind DK, OpenACCGangKind GK,
+                           Expr *E);
 
   // Does the checking for a 'gang' clause that needs to be done in dependent
   // and not dependent cases.
@@ -771,7 +778,7 @@ class SemaOpenACC : public SemaBase {
     SemaOpenACC &SemaRef;
     ComputeConstructInfo OldActiveComputeConstructInfo;
     OpenACCDirectiveKind DirKind;
-    SourceLocation OldLoopGangClauseOnKernelLoc;
+    LoopGangOnKernelTy OldLoopGangClauseOnKernel;
     SourceLocation OldLoopWorkerClauseLoc;
     SourceLocation OldLoopVectorClauseLoc;
     LoopWithoutSeqCheckingInfo OldLoopWithoutSeqInfo;

diff  --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 30491d6f3a3bd6..16348d0b8837ef 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -737,6 +737,24 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
       return nullptr;
     }
   }
+
+  // OpenACC 3.3 Section 2.9.2:
+  // An argument with no keyword or with the 'num' wkeyword is allowed only when
+  // the 'num_gangs' does not appear on the 'kernel' construct.
+  if (Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop) {
+    auto GangClauses = llvm::make_filter_range(
+        ExistingClauses, llvm::IsaPred<OpenACCGangClause>);
+
+    for (auto *GC : GangClauses) {
+      if (cast<OpenACCGangClause>(GC)->hasExprOfKind(OpenACCGangKind::Num)) {
+        SemaRef.Diag(Clause.getBeginLoc(),
+                     diag::err_acc_num_arg_conflict_reverse);
+        SemaRef.Diag(GC->getBeginLoc(), diag::note_acc_previous_clause_here);
+        return nullptr;
+      }
+    }
+  }
+
   return OpenACCNumGangsClause::Create(
       Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getIntExprs(),
       Clause.getEndLoc());
@@ -1033,6 +1051,136 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause(
                                           Clause.getEndLoc());
 }
 
+ExprResult CheckGangStaticExpr(SemaOpenACC &S, Expr *E) {
+  if (isa<OpenACCAsteriskSizeExpr>(E))
+    return E;
+  return S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang,
+                        E->getBeginLoc(), E);
+}
+
+bool IsOrphanLoop(OpenACCDirectiveKind DK, OpenACCDirectiveKind AssocKind) {
+  return DK == OpenACCDirectiveKind::Loop &&
+         AssocKind == OpenACCDirectiveKind::Invalid;
+}
+
+bool HasAssocKind(OpenACCDirectiveKind DK, OpenACCDirectiveKind AssocKind) {
+  return DK == OpenACCDirectiveKind::Loop &&
+         AssocKind != OpenACCDirectiveKind::Invalid;
+}
+
+ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, OpenACCGangKind GK,
+                             OpenACCClauseKind CK, OpenACCDirectiveKind DK,
+                             OpenACCDirectiveKind AssocKind) {
+  S.Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
+      << GK << CK << IsOrphanLoop(DK, AssocKind) << DK
+      << HasAssocKind(DK, AssocKind) << AssocKind;
+  return ExprError();
+}
+
+ExprResult CheckGangParallelExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
+                                 OpenACCDirectiveKind AssocKind,
+                                 OpenACCGangKind GK, Expr *E) {
+  switch (GK) {
+  case OpenACCGangKind::Static:
+    return CheckGangStaticExpr(S, E);
+  case OpenACCGangKind::Num:
+    // 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.
+    return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind);
+  case OpenACCGangKind::Dim: {
+    // 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.
+    if (!E)
+      return ExprError();
+    ExprResult Res =
+        S.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(S.getASTContext());
+
+    if (!ICE || *ICE <= 0 || ICE > 3) {
+      S.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(S.getASTContext(), Res.get(), APValue{*ICE})};
+  }
+  }
+  llvm_unreachable("Unknown gang kind in gang parallel check");
+}
+
+ExprResult CheckGangKernelsExpr(SemaOpenACC &S,
+                                ArrayRef<const OpenACCClause *> ExistingClauses,
+                                OpenACCDirectiveKind DK,
+                                OpenACCDirectiveKind AssocKind,
+                                OpenACCGangKind GK, Expr *E) {
+  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:
+    return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind);
+  case OpenACCGangKind::Num: {
+    // 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.
+
+    // If this is a 'combined' construct, search the list of existing clauses.
+    // Else we need to search the containing 'kernel'.
+    auto Collection = isOpenACCCombinedDirectiveKind(DK)
+                          ? ExistingClauses
+                          : S.getActiveComputeConstructInfo().Clauses;
+
+    const auto *Itr =
+        llvm::find_if(Collection, llvm::IsaPred<OpenACCNumGangsClause>);
+
+    if (Itr != Collection.end()) {
+      S.Diag(E->getBeginLoc(), diag::err_acc_num_arg_conflict)
+          << OpenACCClauseKind::Gang << DK << HasAssocKind(DK, AssocKind)
+          << AssocKind << OpenACCClauseKind::NumGangs;
+
+      S.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+      return ExprError();
+    }
+    return ExprResult{E};
+  }
+  case OpenACCGangKind::Static:
+    return CheckGangStaticExpr(S, E);
+    return ExprError();
+  }
+  llvm_unreachable("Unknown gang kind in gang kernels check");
+}
+
+ExprResult CheckGangSerialExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
+                               OpenACCDirectiveKind AssocKind,
+                               OpenACCGangKind GK, Expr *E) {
+  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:
+    return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind);
+  case OpenACCGangKind::Static:
+    return CheckGangStaticExpr(S, E);
+  }
+  llvm_unreachable("Unknown gang kind in gang serial check");
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   if (DiagIfSeqClause(Clause))
@@ -1054,8 +1202,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
     case OpenACCDirectiveKind::Serial:
       // GCC disallows this, and there is no real good reason for us to permit
       // it, so disallow until we come up with a use case that makes sense.
-      SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
-          << OpenACCClauseKind::Vector << "num" << /*serial=*/3;
+      DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
+                        OpenACCClauseKind::Vector, Clause.getDirectiveKind(),
+                        SemaRef.getActiveComputeConstructInfo().Kind);
       IntExpr = nullptr;
       break;
     case OpenACCDirectiveKind::Kernels: {
@@ -1064,7 +1213,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
                         llvm::IsaPred<OpenACCVectorLengthClause>);
       if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
         SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
-            << OpenACCClauseKind::Vector << /*vector_length=*/2;
+            << OpenACCClauseKind::Vector << Clause.getDirectiveKind()
+            << HasAssocKind(Clause.getDirectiveKind(),
+                            SemaRef.getActiveComputeConstructInfo().Kind)
+            << SemaRef.getActiveComputeConstructInfo().Kind
+            << OpenACCClauseKind::VectorLength;
         SemaRef.Diag((*Itr)->getBeginLoc(),
                      diag::note_acc_previous_clause_here);
 
@@ -1113,18 +1266,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
   if (IntExpr) {
     switch (SemaRef.getActiveComputeConstructInfo().Kind) {
     case OpenACCDirectiveKind::Invalid:
-      SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
-          << OpenACCClauseKind::Worker << "num" << /*orphan=*/0;
-      IntExpr = nullptr;
-      break;
     case OpenACCDirectiveKind::Parallel:
-      SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
-          << OpenACCClauseKind::Worker << "num" << /*parallel=*/1;
-      IntExpr = nullptr;
-      break;
     case OpenACCDirectiveKind::Serial:
-      SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
-          << OpenACCClauseKind::Worker << "num" << /*serial=*/3;
+      DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
+                        OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
+                        SemaRef.getActiveComputeConstructInfo().Kind);
       IntExpr = nullptr;
       break;
     case OpenACCDirectiveKind::Kernels: {
@@ -1133,7 +1279,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
                         llvm::IsaPred<OpenACCNumWorkersClause>);
       if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
         SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
-            << OpenACCClauseKind::Worker << /*num_workers=*/1;
+            << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
+            << HasAssocKind(Clause.getDirectiveKind(),
+                            SemaRef.getActiveComputeConstructInfo().Kind)
+            << SemaRef.getActiveComputeConstructInfo().Kind
+            << OpenACCClauseKind::NumWorkers;
         SemaRef.Diag((*Itr)->getBeginLoc(),
                      diag::note_acc_previous_clause_here);
 
@@ -1187,12 +1337,15 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
   // 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)
+  if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop &&
+      !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
     return isNotImplemented();
 
   // OpenACC 3.3 Section 2.9.11: A reduction clause may not appear on a loop
   // directive that has a gang clause and is within a compute construct that has
   // a num_gangs clause with more than one explicit argument.
+  // TODO OpenACC: When we implement reduction on combined constructs, we need
+  // to do this too.
   if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop &&
       SemaRef.getActiveComputeConstructInfo().Kind !=
           OpenACCDirectiveKind::Invalid) {
@@ -1229,31 +1382,13 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
 
   for (unsigned I = 0; I < Clause.getIntExprs().size(); ++I) {
     OpenACCGangKind GK = Clause.getGangKinds()[I];
-    ExprResult ER = SemaRef.CheckGangExpr(GK, Clause.getIntExprs()[I]);
+    ExprResult ER =
+        SemaRef.CheckGangExpr(ExistingClauses, Clause.getDirectiveKind(), 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_num_arg_conflict)
-            << OpenACCClauseKind::Gang << /*num_gangs=*/0;
-        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()) {
@@ -1269,47 +1404,50 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
     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_clause_in_clause_region)
-        << OpenACCClauseKind::Gang << OpenACCClauseKind::Gang
-        << /*kernels construct info*/ 1;
-    SemaRef.Diag(SemaRef.LoopGangClauseOnKernelLoc,
-                 diag::note_acc_previous_clause_here);
-    return nullptr;
-  }
+  if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+    // 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.LoopGangClauseOnKernel.Loc.isValid()) {
+      // This handles the 'inner loop' diagnostic, but we cannot set that we're
+      // on one of these until we get to the end of the construct.
+      SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+          << OpenACCClauseKind::Gang << OpenACCClauseKind::Gang
+          << /*kernels construct info*/ 1
+          << SemaRef.LoopGangClauseOnKernel.DirKind;
+      SemaRef.Diag(SemaRef.LoopGangClauseOnKernel.Loc,
+                   diag::note_acc_previous_clause_here);
+      return nullptr;
+    }
 
-  // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not
-  // contain a loop with a gang or worker clause unless within a nested compute
-  // region.
-  if (SemaRef.LoopWorkerClauseLoc.isValid()) {
-    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
-    // one of these until we get to the end of the construct.
-    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
-        << OpenACCClauseKind::Gang << OpenACCClauseKind::Worker
-        << /*kernels construct info*/ 1;
-    SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
-                 diag::note_acc_previous_clause_here);
-    return nullptr;
-  }
+    // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not
+    // contain a loop with a gang or worker clause unless within a nested
+    // compute region.
+    if (SemaRef.LoopWorkerClauseLoc.isValid()) {
+      // This handles the 'inner loop' diagnostic, but we cannot set that we're
+      // on one of these until we get to the end of the construct.
+      SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+          << OpenACCClauseKind::Gang << OpenACCClauseKind::Worker
+          << /*!kernels construct info*/ 0;
+      SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
+                   diag::note_acc_previous_clause_here);
+      return nullptr;
+    }
 
-  // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
-  // contain a loop with a gang, worker, or vector clause unless within a nested
-  // compute region.
-  if (SemaRef.LoopVectorClauseLoc.isValid()) {
-    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
-    // one of these until we get to the end of the construct.
-    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
-        << OpenACCClauseKind::Gang << OpenACCClauseKind::Vector
-        << /*kernels construct info*/ 1;
-    SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
-                 diag::note_acc_previous_clause_here);
-    return nullptr;
+    // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
+    // contain a loop with a gang, worker, or vector clause unless within a
+    // nested compute region.
+    if (SemaRef.LoopVectorClauseLoc.isValid()) {
+      // This handles the 'inner loop' diagnostic, but we cannot set that we're
+      // on one of these until we get to the end of the construct.
+      SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+          << OpenACCClauseKind::Gang << OpenACCClauseKind::Vector
+          << /*!kernels construct info*/ 0;
+      SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+                   diag::note_acc_previous_clause_here);
+      return nullptr;
+    }
   }
 
   return SemaRef.CheckGangClause(ExistingClauses, Clause.getBeginLoc(),
@@ -1489,7 +1627,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     ArrayRef<const OpenACCClause *> UnInstClauses,
     ArrayRef<OpenACCClause *> Clauses)
     : SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo),
-      DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc),
+      DirKind(DK), OldLoopGangClauseOnKernel(S.LoopGangClauseOnKernel),
       OldLoopWorkerClauseLoc(S.LoopWorkerClauseLoc),
       OldLoopVectorClauseLoc(S.LoopVectorClauseLoc),
       OldLoopWithoutSeqInfo(S.LoopWithoutSeqInfo),
@@ -1510,7 +1648,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     // within a nested compute region.
     //
     // Implement the 'unless within a nested compute region' part.
-    SemaRef.LoopGangClauseOnKernelLoc = {};
+    SemaRef.LoopGangClauseOnKernel = {};
     SemaRef.LoopWorkerClauseLoc = {};
     SemaRef.LoopVectorClauseLoc = {};
     SemaRef.LoopWithoutSeqInfo = {};
@@ -1524,8 +1662,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
     SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
 
-    // TODO: OpenACC: We need to set these 3, CollapseInfo, and TileInfo
-    SemaRef.LoopGangClauseOnKernelLoc = {};
+    SemaRef.LoopGangClauseOnKernel = {};
     SemaRef.LoopWorkerClauseLoc = {};
     SemaRef.LoopVectorClauseLoc = {};
 
@@ -1548,7 +1685,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
       // 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();
+        SemaRef.LoopGangClauseOnKernel = {(*Itr)->getBeginLoc(), DirKind};
     }
 
     if (UnInstClauses.empty()) {
@@ -1586,7 +1723,8 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
       // 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();
+        SemaRef.LoopGangClauseOnKernel = {(*Itr)->getBeginLoc(),
+                                          OpenACCDirectiveKind::Kernels};
     }
 
     if (UnInstClauses.empty()) {
@@ -1671,7 +1809,7 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
 
 SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
   SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
-  SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc;
+  SemaRef.LoopGangClauseOnKernel = OldLoopGangClauseOnKernel;
   SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc;
   SemaRef.LoopVectorClauseLoc = OldLoopVectorClauseLoc;
   SemaRef.LoopWithoutSeqInfo = OldLoopWithoutSeqInfo;
@@ -2328,106 +2466,47 @@ 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_int_arg_invalid)
-          << OpenACCClauseKind::Gang << 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_int_arg_invalid)
-          << OpenACCClauseKind::Gang << GK << /*kernels=*/2;
-      return ExprError();
-    // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
-    // construct, the gang clause behaves as follows. ... An argument with no
-    // 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_int_arg_invalid)
-          << OpenACCClauseKind::Gang << GK << /*Kernels=*/3;
-      return ExprError();
-    case OpenACCGangKind::Static:
-      return CheckGangStaticExpr(*this, E);
+ExprResult
+SemaOpenACC::CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses,
+                           OpenACCDirectiveKind DK, OpenACCGangKind GK,
+                           Expr *E) {
+  // There are two cases for the enforcement here: the 'current' directive is a
+  // 'loop', where we need to check the active compute construct kind, or the
+  // current directive is a 'combined' construct, where we have to check the
+  // current one.
+  switch (DK) {
+  case OpenACCDirectiveKind::ParallelLoop:
+    return CheckGangParallelExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK,
+                                 E);
+  case OpenACCDirectiveKind::SerialLoop:
+    return CheckGangSerialExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK,
+                               E);
+  case OpenACCDirectiveKind::KernelsLoop:
+    return CheckGangKernelsExpr(*this, ExistingClauses, DK,
+                                ActiveComputeConstructInfo.Kind, GK, E);
+  case OpenACCDirectiveKind::Loop:
+    switch (ActiveComputeConstructInfo.Kind) {
+    case OpenACCDirectiveKind::Invalid:
+    case OpenACCDirectiveKind::Parallel:
+    case OpenACCDirectiveKind::ParallelLoop:
+      return CheckGangParallelExpr(*this, DK, ActiveComputeConstructInfo.Kind,
+                                   GK, E);
+    case OpenACCDirectiveKind::SerialLoop:
+    case OpenACCDirectiveKind::Serial:
+      return CheckGangSerialExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK,
+                                 E);
+    case OpenACCDirectiveKind::KernelsLoop:
+    case OpenACCDirectiveKind::Kernels:
+      return CheckGangKernelsExpr(*this, ExistingClauses, DK,
+                                  ActiveComputeConstructInfo.Kind, GK, E);
+    default:
+      llvm_unreachable("Non compute construct in active compute construct?");
     }
-  } break;
   default:
-    llvm_unreachable("Non compute construct in active compute construct?");
+    // TODO: OpenACC: when we implement this on 'routine', we'll have to
+    // implement its checking here.
+    llvm_unreachable("Invalid directive kind for a Gang clause");
   }
-
   llvm_unreachable("Compute construct directive not handled?");
 }
 

diff  --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 3a8f2d95f329b8..81e515e7cb2a9a 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -12027,7 +12027,9 @@ void OpenACCClauseTransform<Derived>::VisitGangClause(
     if (!ER.isUsable())
       continue;
 
-    ER = Self.getSema().OpenACC().CheckGangExpr(C.getExpr(I).first, ER.get());
+    ER = Self.getSema().OpenACC().CheckGangExpr(ExistingClauses,
+                                                ParsedClause.getDirectiveKind(),
+                                                C.getExpr(I).first, ER.get());
     if (!ER.isUsable())
       continue;
     TransformedGangKinds.push_back(C.getExpr(I).first);

diff  --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp
index 435c770c7457d1..40f174e539c01e 100644
--- a/clang/test/AST/ast-print-openacc-combined-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp
@@ -242,4 +242,75 @@ void foo() {
 #pragma acc parallel loop vector_length((int)array[1])
   for(int i = 0;i<5;++i);
 
+// CHECK: #pragma acc parallel loop gang(dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(dim:2)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc serial loop gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop gang(static:i)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i) gang(dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i) gang(dim:2)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i, dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i, dim:2)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(dim:2)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i) gang(dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i) gang(dim:2)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i, dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i, dim:2)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc kernels loop gang(num: i) gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop gang(i) gang(static:i)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc kernels loop gang(num: i) gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop gang(num:i) gang(static:i)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc serial loop gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop gang(static:i)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc serial loop gang(static: *)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop gang(static:*)
+  for(int i = 0;i<5;++i);
 }

diff  --git a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
index a6f57a63a91ddf..e35bd6da2f18b3 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -155,7 +155,6 @@ void uses() {
 #pragma acc parallel loop auto tile(1+2, 1)
   for(unsigned j = 0; j < 5; ++j)
     for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc parallel loop auto gang
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto wait
@@ -276,7 +275,6 @@ void uses() {
 #pragma acc parallel loop tile(1+2, 1) auto
   for(unsigned j = 0; j < 5; ++j)
     for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc parallel loop gang auto
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop wait auto
@@ -398,7 +396,6 @@ void uses() {
 #pragma acc parallel loop independent tile(1+2, 1)
   for(unsigned j = 0; j < 5; ++j)
     for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc parallel loop independent gang
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent wait
@@ -519,7 +516,6 @@ void uses() {
 #pragma acc parallel loop tile(1+2, 1) independent
   for(unsigned j = 0; j < 5; ++j)
     for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc parallel loop gang independent
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop wait independent
@@ -650,9 +646,8 @@ void uses() {
 #pragma acc parallel loop seq wait
   for(unsigned i = 0; i < 5; ++i);
 
-  // TODOexpected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'parallel loop' construct}}
-  // TODOexpected-note at +1{{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 'parallel loop' construct}}
+  // expected-note at +1{{previous clause is here}}
 #pragma acc parallel loop gang seq
   for(unsigned i = 0; i < 5; ++i);
   // TODOexpected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'parallel loop' construct}}

diff  --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index 9a60fb4c665e5a..d3ed8234d16b14 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -212,7 +212,6 @@ void uses() {
   for(int j = 0; j < 5; ++j)
     for(int i = 0; i < 5; ++i);
 
-  // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc serial loop dtype(*) gang
   for(int i = 0; i < 5; ++i);
 #pragma acc parallel loop device_type(*) wait

diff  --git a/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp b/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp
new file mode 100644
index 00000000000000..f179b928215e71
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp
@@ -0,0 +1,216 @@
+
+// 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
+
+  int Val;
+
+#pragma acc parallel loop gang(dim:1)
+  for(int i = 0; i < 5; ++i);
+  // CHECK: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause dim
+  // CHECK-NEXT: ConstantExpr
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop gang(static:Val)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc kernels loop gang(num:1) gang(static:Val)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: gang clause num
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop gang(dim:1, static:Val)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause dim static
+  // CHECK-NEXT: ConstantExpr
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc serial loop gang(static:Val)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc serial loop gang(static:*)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: OpenACCAsteriskSizeExpr
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc serial loop gang
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+  // CHECK-NEXT: gang clause
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc kernels loop gang(num:1)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: gang clause num
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+}
+
+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
+
+#pragma acc parallel loop gang(dim:One)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause dim
+  // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop gang(static:Val)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+
+#pragma acc parallel loop gang(static:*)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: OpenACCAsteriskSizeExpr
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop gang(dim:One) gang(static:Val)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause dim
+  // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int'
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop gang(dim:One, static:Val)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause dim static
+  // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc serial loop gang
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+  // CHECK-NEXT: gang clause
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // 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: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause dim
+  // CHECK-NEXT: ConstantExpr
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr
+  // CHECK-NEXT: NonTypeTemplateParmDecl
+  // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: OpenACCAsteriskSizeExpr
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause dim
+  // CHECK-NEXT: ConstantExpr
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr
+  // CHECK-NEXT: NonTypeTemplateParmDecl
+  // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+  // CHECK-NEXT: gang clause static
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: gang clause dim static
+  // CHECK-NEXT: ConstantExpr
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr
+  // CHECK-NEXT: NonTypeTemplateParmDecl
+  // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+  // CHECK-NEXT: gang clause
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+}
+
+void inst() {
+  TemplateUses<int, 1>(5);
+}
+
+#endif // PCH_HELPER

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

diff  --git a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
index a2bda5a7e82f63..6b4c03bbc50c6b 100644
--- a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
+++ b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
@@ -278,7 +278,7 @@ void Kernels() {
 #pragma acc kernels
 #pragma acc loop gang(num:1)
   for(int j = 0; j < 5; ++j) {
-    // 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-error at +2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels' construct}}
     // expected-note at -3{{previous clause is here}}
 #pragma acc loop gang(static:1)
     for(int i = 0; i < 5; ++i);


        


More information about the cfe-commits mailing list