[OpenMP][Offload] Add support for cgroup modifier in dyn_groupprivate

Author: kevinsala

clang/include/clang/AST/OpenMPClause.h

@@ -9772,63 +9772,155 @@ class OMPXDynCGroupMemClause
 /// and '#pragma omp teams ...' directives.
 ///
 /// \code
-/// #pragma omp target [...] dyn_groupprivate(N)
+/// #pragma omp target [...] dyn_groupprivate(a,b: N)
 /// \endcode
-class OMPDynGroupprivateClause
-    : public OMPOneStmtClause<llvm::omp::OMPC_dyn_groupprivate, OMPClause>,
-      public OMPClauseWithPreInit {
+class OMPDynGroupprivateClause : public OMPClause, public OMPClauseWithPreInit {
   friend class OMPClauseReader;
 
-  /// Modifiers for 'grainsize' clause.
-  OpenMPDynGroupprivateClauseModifier Modifier = OMPC_DYN_GROUPPRIVATE_unknown;
+  /// Location of '('.
+  SourceLocation LParenLoc;
 
-  /// Location of the modifier.
-  SourceLocation ModifierLoc;
+  /// Modifiers for 'dyn_groupprivate' clause.
+  enum {FIRST, SECOND, NUM_MODIFIERS};
+  OpenMPDynGroupprivateClauseModifier Modifiers[NUM_MODIFIERS];
 
-  /// Set size.
-  void setSize(Expr *E) { setStmt(E); }
+  /// Locations of modifiers.
+  SourceLocation ModifiersLoc[NUM_MODIFIERS];
 
-  /// Sets modifier.
-  void setModifier(OpenMPDynGroupprivateClauseModifier M) { Modifier = M; }
+  /// The size of the dyn_groupprivate.
+  Expr *Size = nullptr;
 
-  /// Sets modifier location.
-  void setModifierLoc(SourceLocation Loc) { ModifierLoc = Loc; }
+  /// Set the first dyn_groupprivate modifier.
+  ///
+  /// \param M The modifier.
+  void setFirstDynGroupprivateModifier(OpenMPDynGroupprivateClauseModifier M) {
+    Modifiers[FIRST] = M;
+  }
+
+  /// Set the second dyn_groupprivate modifier.
+  ///
+  /// \param M The modifier.
+  void setSecondDynGroupprivateModifier(OpenMPDynGroupprivateClauseModifier M) {
+    Modifiers[SECOND] = M;
+  }
+
+  /// Set location of the first dyn_groupprivate modifier.
+  void setFirstDynGroupprivateModifierLoc(SourceLocation Loc) {
+    ModifiersLoc[FIRST] = Loc;
+  }
+
+  /// Set location of the second dyn_groupprivate modifier.
+  void setSecondDynGroupprivateModifierLoc(SourceLocation Loc) {
+    ModifiersLoc[SECOND] = Loc;
+  }
+
+  /// Set dyn_groupprivate modifier location.
+  ///
+  /// \param M The modifier location.
+  void setDynGroupprivateModifer(OpenMPDynGroupprivateClauseModifier M) {
+    if (Modifiers[FIRST] == OMPC_DYN_GROUPPRIVATE_unknown)
+      Modifiers[FIRST] = M;
+    else {
+      assert(Modifiers[SECOND] == OMPC_DYN_GROUPPRIVATE_unknown);
+      Modifiers[SECOND] = M;
+    }
+  }
+
+  /// Sets the location of '('.
+  ///
+  /// \param Loc Location of '('.
+  void setLParenLoc(SourceLocation Loc) { LParenLoc = Loc; }
+
+  /// Set size.
+  ///
+  /// \param E Size.
+  void setSize(Expr *E) { Size = E; }
 
 public:
-  /// Build 'dyn_groupprivate' clause.
+  /// Build 'dyn_groupprivate' clause with a size expression \a Size.
   ///
-  /// \param Modifier Clause modifier.
-  /// \param Size Size expression.
-  /// \param HelperSize Helper Size expression
-  /// \param CaptureRegion Innermost OpenMP region where expressions in this
   /// \param StartLoc Starting location of the clause.
   /// \param LParenLoc Location of '('.
-  /// \param ModifierLoc Modifier location.
   /// \param EndLoc Ending location of the clause.
-  OMPDynGroupprivateClause(OpenMPDynGroupprivateClauseModifier Modifier,
-                         Expr *Size, Stmt *HelperSize,
-                         OpenMPDirectiveKind CaptureRegion,
-                         SourceLocation StartLoc, SourceLocation LParenLoc,
-                         SourceLocation ModifierLoc, SourceLocation EndLoc)
-      : OMPOneStmtClause(Size, StartLoc, LParenLoc, EndLoc),
-        OMPClauseWithPreInit(this), Modifier(Modifier), ModifierLoc(ModifierLoc) {
+  /// \param Size Size.
+  /// \param M1 The first modifier applied to 'dyn_groupprivate' clause.
+  /// \param M1Loc Location of the first modifier.
+  /// \param M2 The second modifier applied to 'dyn_groupprivate' clause.
+  /// \param M2Loc Location of the second modifier.
+  OMPDynGroupprivateClause(SourceLocation StartLoc, SourceLocation LParenLoc,
+                    SourceLocation EndLoc,
+                    Expr *Size, Stmt *HelperSize,
+                    OpenMPDirectiveKind CaptureRegion,
+                    OpenMPDynGroupprivateClauseModifier M1,
+                    SourceLocation M1Loc,
+                    OpenMPDynGroupprivateClauseModifier M2,
+                    SourceLocation M2Loc)
+      : OMPClause(llvm::omp::OMPC_dyn_groupprivate, StartLoc, EndLoc),
+        OMPClauseWithPreInit(this), LParenLoc(LParenLoc), Size(Size) {
     setPreInitStmt(HelperSize, CaptureRegion);
+    Modifiers[FIRST] = M1;
+    Modifiers[SECOND] = M2;
+    ModifiersLoc[FIRST] = M1Loc;
+    ModifiersLoc[SECOND] = M2Loc;
   }
 
   /// Build an empty clause.
-  OMPDynGroupprivateClause() : OMPOneStmtClause(), OMPClauseWithPreInit(this) {}
+  explicit OMPDynGroupprivateClause()
+      : OMPClause(llvm::omp::OMPC_dyn_groupprivate, SourceLocation(), SourceLocation()),
+        OMPClauseWithPreInit(this) {
+    Modifiers[FIRST] = OMPC_DYN_GROUPPRIVATE_unknown;
+    Modifiers[SECOND] = OMPC_DYN_GROUPPRIVATE_unknown;
+  }
 
-  /// Return the size expression.
-  Expr *getSize() { return getStmtAs<Expr>(); }
+  /// Get the first modifier of the clause.
+  OpenMPDynGroupprivateClauseModifier getFirstDynGroupprivateModifier() const {
+    return Modifiers[FIRST];
+  }
 
-  /// Return the size expression.
-  Expr *getSize() const { return getStmtAs<Expr>(); }
+  /// Get the second modifier of the clause.
+  OpenMPDynGroupprivateClauseModifier getSecondDynGroupprivateModifier() const {
+    return Modifiers[SECOND];
+  }
 
-  /// Gets modifier.
-  OpenMPDynGroupprivateClauseModifier getModifier() const { return Modifier; }
+  /// Get location of '('.
+  SourceLocation getLParenLoc() { return LParenLoc; }
 
-  /// Gets modifier location.
-  SourceLocation getModifierLoc() const { return ModifierLoc; }
+  /// Get the first modifier location.
+  SourceLocation getFirstDynGroupprivateModifierLoc() const {
+    return ModifiersLoc[FIRST];
+  }
+
+  /// Get the second modifier location.
+  SourceLocation getSecondDynGroupprivateModifierLoc() const {
+    return ModifiersLoc[SECOND];
+  }
+
+  /// Get size.
+  Expr *getSize() { return Size; }
+
+  /// Get size.
+  const Expr *getSize() const { return Size; }
+
+  child_range children() {
+    return child_range(reinterpret_cast<Stmt **>(&Size),
+                       reinterpret_cast<Stmt **>(&Size) + 1);
+  }
+
+  const_child_range children() const {
+    auto Children = const_cast<OMPDynGroupprivateClause *>(this)->children();
+    return const_child_range(Children.begin(), Children.end());
+  }
+
+  child_range used_children() {
+    return child_range(child_iterator(), child_iterator());
+  }
+  const_child_range used_children() const {
+    return const_child_range(const_child_iterator(), const_child_iterator());
+  }
+
+  static bool classof(const OMPClause *T) {
+    return T->getClauseKind() == llvm::omp::OMPC_dyn_groupprivate;
+  }
 };
 
 /// This represents the 'doacross' clause for the '#pragma omp ordered'

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11995,6 +11995,8 @@ def err_omp_unexpected_schedule_modifier : Error<
   "modifier '%0' cannot be used along with modifier '%1'">;
 def err_omp_schedule_nonmonotonic_static : Error<
   "'nonmonotonic' modifier can only be specified with 'dynamic' or 'guided' schedule kind">;
+def err_omp_unexpected_dyn_groupprivate_modifier : Error<
+  "modifier '%0' cannot be used along with modifier '%1' in dyn_groupprivate">;
 def err_omp_simple_clause_incompatible_with_ordered : Error<
   "'%0' clause with '%1' modifier cannot be specified if an 'ordered' clause is specified">;
 def err_omp_ordered_simd : Error<

clang/include/clang/Basic/OpenMPKinds.def

@@ -231,6 +231,7 @@ OPENMP_BIND_KIND(thread)
 OPENMP_GRAINSIZE_MODIFIER(strict)
 
 // Modifiers for the 'dyn_groupprivate' clause.
+OPENMP_DYN_GROUPPRIVATE_MODIFIER(cgroup)
 OPENMP_DYN_GROUPPRIVATE_MODIFIER(strict)
 OPENMP_DYN_GROUPPRIVATE_MODIFIER(fallback)
 

clang/include/clang/Basic/OpenMPKinds.h

@@ -223,6 +223,10 @@ enum OpenMPDynGroupprivateClauseModifier {
   OMPC_DYN_GROUPPRIVATE_unknown
 };
 
+/// Number of allowed dyn_groupprivate-modifiers.
+static constexpr unsigned NumberOfOMPDynGroupprivateClauseModifiers =
+    OMPC_DYN_GROUPPRIVATE_unknown;
+
 enum OpenMPNumTasksClauseModifier {
 #define OPENMP_NUMTASKS_MODIFIER(Name) OMPC_NUMTASKS_##Name,
 #include "clang/Basic/OpenMPKinds.def"

clang/include/clang/Sema/SemaOpenMP.h

@@ -1386,11 +1386,10 @@ class SemaOpenMP : public SemaBase {
                                             SourceLocation EndLoc);
 
   /// Called on a well-formed 'dyn_groupprivate' clause.
-  OMPClause *
-  ActOnOpenMPDynGroupprivateClause(OpenMPDynGroupprivateClauseModifier Modifier,
-                                       Expr *Size, SourceLocation StartLoc,
-                                       SourceLocation LParenLoc,
-                                       SourceLocation ModifierLoc, SourceLocation EndLoc);
+  OMPClause *ActOnOpenMPDynGroupprivateClause(
+      OpenMPDynGroupprivateClauseModifier M1, OpenMPDynGroupprivateClauseModifier M2,
+      Expr *Size, SourceLocation StartLoc, SourceLocation LParenLoc,
+      SourceLocation M1Loc, SourceLocation M2Loc, SourceLocation EndLoc);
 
   /// Called on well-formed 'doacross' clause.
   OMPClause *

clang/lib/AST/OpenMPClause.cpp

@@ -2727,13 +2727,17 @@ void OMPClausePrinter::VisitOMPXDynCGroupMemClause(
   OS << ")";
 }
 
-void OMPClausePrinter::VisitOMPDynGroupprivateClause(
-    OMPDynGroupprivateClause *Node) {
+void OMPClausePrinter::VisitOMPDynGroupprivateClause(OMPDynGroupprivateClause *Node) {
   OS << "dyn_groupprivate(";
-  OpenMPDynGroupprivateClauseModifier Modifier = Node->getModifier();
-  if (Modifier != OMPC_DYN_GROUPPRIVATE_unknown) {
-    OS << getOpenMPSimpleClauseTypeName(Node->getClauseKind(), Modifier)
-       << ": ";
+  if (Node->getFirstDynGroupprivateModifier() != OMPC_SCHEDULE_MODIFIER_unknown) {
+    OS << getOpenMPSimpleClauseTypeName(OMPC_dyn_groupprivate,
+                                        Node->getFirstDynGroupprivateModifier());
+    if (Node->getSecondDynGroupprivateModifier() != OMPC_SCHEDULE_MODIFIER_unknown) {
+      OS << ", ";
+      OS << getOpenMPSimpleClauseTypeName(OMPC_dyn_groupprivate,
+                                          Node->getSecondDynGroupprivateModifier());
+    }
+    OS << ": ";
   }
   Node->getSize()->printPretty(OS, nullptr, Policy, 0);
   OS << ")";

clang/lib/AST/StmtProfile.cpp

@@ -960,7 +960,7 @@ void OMPClauseProfiler::VisitOMPXDynCGroupMemClause(
 void OMPClauseProfiler::VisitOMPDynGroupprivateClause(
     const OMPDynGroupprivateClause *C) {
   VistOMPClauseWithPreInit(C);
-  if (Expr *Size = C->getSize())
+  if (auto *Size = C->getSize())
     Profiler->VisitStmt(Size);
 }
 void OMPClauseProfiler::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -9500,8 +9500,8 @@ static std::pair<llvm::Value *, bool> emitDynCGroupMem(const OMPExecutableDirect
         DynGPClause->getSize(), /*IgnoreResultAssign=*/true);
     DynGP = CGF.Builder.CreateIntCast(DynGPVal, CGF.Int32Ty,
                                       /*isSigned=*/false);
-    DynGPFallback = (DynGPClause->getModifier() == OMPC_DYN_GROUPPRIVATE_fallback ||
-        OMPC_DYN_GROUPPRIVATE_unknown);
+    DynGPFallback = (DynGPClause->getFirstDynGroupprivateModifier() != OMPC_DYN_GROUPPRIVATE_strict &&
+        DynGPClause->getSecondDynGroupprivateModifier() != OMPC_DYN_GROUPPRIVATE_strict);
   } else if (auto *OMPXDynCGClause = D.getSingleClause<OMPXDynCGroupMemClause>()) {
     CodeGenFunction::RunCleanupsScope DynCGMemScope(CGF);
     llvm::Value *DynCGMemVal = CGF.EmitScalarExpr(

clang/lib/Parse/ParseOpenMP.cpp

@@ -3837,35 +3837,37 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
       KLoc.emplace_back();
     }
   } else if (Kind == OMPC_dyn_groupprivate) {
-    // Parse optional <dyn_groupprivate modifier> ':'
-    OpenMPDynGroupprivateClauseModifier Modifier =
-        static_cast<OpenMPDynGroupprivateClauseModifier>(getOpenMPSimpleClauseType(
-            Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
-            getLangOpts()));
-    if (getLangOpts().OpenMP >= 51) {
-      if (NextToken().is(tok::colon)) {
-        Arg.push_back(Modifier);
-        KLoc.push_back(Tok.getLocation());
-        // Parse modifier
+    enum { Modifier1, Modifier2, NumberOfElements };
+    Arg.resize(NumberOfElements);
+    KLoc.resize(NumberOfElements);
+    Arg[Modifier1] = OMPC_DYN_GROUPPRIVATE_unknown;
+    Arg[Modifier2] = OMPC_DYN_GROUPPRIVATE_unknown;
+    unsigned Modifier = getOpenMPSimpleClauseType(
+        Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts());
+
+    if (Modifier < OMPC_DYN_GROUPPRIVATE_unknown) {
+      // Parse 'modifier'
+      Arg[Modifier1] = Modifier;
+      KLoc[Modifier1] = Tok.getLocation();
+      if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
+          Tok.isNot(tok::annot_pragma_openmp_end))
         ConsumeAnyToken();
-        // Parse ':'
+      if (Tok.is(tok::comma)) {
+        // Parse ',' 'modifier'
         ConsumeAnyToken();
-      } else {
-        if (Modifier == OMPC_DYN_GROUPPRIVATE_strict) {
-          Diag(Tok, diag::err_modifier_expected_colon) << "strict";
-          // Parse modifier
-          ConsumeAnyToken();
-        } else if (Modifier == OMPC_DYN_GROUPPRIVATE_fallback) {
-          Diag(Tok, diag::err_modifier_expected_colon) << "fallback";
-          // Parse modifier
+        Modifier = getOpenMPSimpleClauseType(
+            Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts());
+        Arg[Modifier2] = Modifier;
+        KLoc[Modifier2] = Tok.getLocation();
+        if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
+            Tok.isNot(tok::annot_pragma_openmp_end))
           ConsumeAnyToken();
-        }
-        Arg.push_back(OMPC_DYN_GROUPPRIVATE_unknown);
-        KLoc.emplace_back();
       }
-    } else {
-      Arg.push_back(OMPC_DYN_GROUPPRIVATE_unknown);
-      KLoc.emplace_back();
+      // Parse ':'
+      if (Tok.is(tok::colon))
+        ConsumeAnyToken();
+      else
+        Diag(Tok, diag::warn_pragma_expected_colon) << "dyn_groupprivate modifier";
     }
   } else if (Kind == OMPC_num_tasks) {
     // Parse optional <num_tasks modifier> ':'