[OpenMP] Add parser/semantic support for dyn_groupprivate clause #152651
+605
−8
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
clang
|
@llvm/pr-subscribers-offload @llvm/pr-subscribers-flang-openmp Author: Kevin Sala Penades (kevinsala) ChangesThis PR adds support for the The implementation reuses part of the runtime support for Patch is 92.16 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/152651.diff 45 Files Affected:
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 1118d3e062e68..2d6931ee052d8 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -9768,6 +9768,161 @@ class OMPXDynCGroupMemClause
Expr *getSize() const { return getStmtAs<Expr>(); }
};
+/// This represents 'dyn_groupprivate' clause in '#pragma omp target ...'
+/// and '#pragma omp teams ...' directives.
+///
+/// \code
+/// #pragma omp target [...] dyn_groupprivate(a,b: N)
+/// \endcode
+class OMPDynGroupprivateClause : public OMPClause, public OMPClauseWithPreInit {
+ friend class OMPClauseReader;
+
+ /// Location of '('.
+ SourceLocation LParenLoc;
+
+ /// Modifiers for 'dyn_groupprivate' clause.
+ enum {FIRST, SECOND, NUM_MODIFIERS};
+ OpenMPDynGroupprivateClauseModifier Modifiers[NUM_MODIFIERS];
+
+ /// Locations of modifiers.
+ SourceLocation ModifiersLoc[NUM_MODIFIERS];
+
+ /// The size of the dyn_groupprivate.
+ Expr *Size = nullptr;
+
+ /// 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 with a size expression \a Size.
+ ///
+ /// \param StartLoc Starting location of the clause.
+ /// \param LParenLoc Location of '('.
+ /// \param EndLoc Ending location of the clause.
+ /// \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.
+ explicit OMPDynGroupprivateClause()
+ : OMPClause(llvm::omp::OMPC_dyn_groupprivate, SourceLocation(), SourceLocation()),
+ OMPClauseWithPreInit(this) {
+ Modifiers[FIRST] = OMPC_DYN_GROUPPRIVATE_unknown;
+ Modifiers[SECOND] = OMPC_DYN_GROUPPRIVATE_unknown;
+ }
+
+ /// Get the first modifier of the clause.
+ OpenMPDynGroupprivateClauseModifier getFirstDynGroupprivateModifier() const {
+ return Modifiers[FIRST];
+ }
+
+ /// Get the second modifier of the clause.
+ OpenMPDynGroupprivateClauseModifier getSecondDynGroupprivateModifier() const {
+ return Modifiers[SECOND];
+ }
+
+ /// Get location of '('.
+ SourceLocation getLParenLoc() { return LParenLoc; }
+
+ /// 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'
/// directive.
///
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 5cb2f57edffe4..129115d56fe82 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -4060,6 +4060,14 @@ bool RecursiveASTVisitor<Derived>::VisitOMPXDynCGroupMemClause(
return true;
}
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPDynGroupprivateClause(
+ OMPDynGroupprivateClause *C) {
+ TRY_TO(VisitOMPClauseWithPreInit(C));
+ TRY_TO(TraverseStmt(C->getSize()));
+ return true;
+}
+
template <typename Derived>
bool RecursiveASTVisitor<Derived>::VisitOMPDoacrossClause(
OMPDoacrossClause *C) {
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index b285309e0b3ca..45b2a79d11561 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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<
diff --git a/clang/include/clang/Basic/OpenMPKinds.def b/clang/include/clang/Basic/OpenMPKinds.def
index 9d6f816eea91f..3321e19cae9b1 100644
--- a/clang/include/clang/Basic/OpenMPKinds.def
+++ b/clang/include/clang/Basic/OpenMPKinds.def
@@ -83,6 +83,9 @@
#ifndef OPENMP_GRAINSIZE_MODIFIER
#define OPENMP_GRAINSIZE_MODIFIER(Name)
#endif
+#ifndef OPENMP_DYN_GROUPPRIVATE_MODIFIER
+#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name)
+#endif
#ifndef OPENMP_NUMTASKS_MODIFIER
#define OPENMP_NUMTASKS_MODIFIER(Name)
#endif
@@ -227,6 +230,11 @@ OPENMP_BIND_KIND(thread)
// Modifiers for the 'grainsize' clause.
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)
+
// Modifiers for the 'num_tasks' clause.
OPENMP_NUMTASKS_MODIFIER(strict)
@@ -245,6 +253,7 @@ OPENMP_DOACROSS_MODIFIER(source_omp_cur_iteration)
#undef OPENMP_NUMTASKS_MODIFIER
#undef OPENMP_NUMTHREADS_MODIFIER
+#undef OPENMP_DYN_GROUPPRIVATE_MODIFIER
#undef OPENMP_GRAINSIZE_MODIFIER
#undef OPENMP_BIND_KIND
#undef OPENMP_ADJUST_ARGS_KIND
diff --git a/clang/include/clang/Basic/OpenMPKinds.h b/clang/include/clang/Basic/OpenMPKinds.h
index f40db4c13c55a..3e164bf1adf22 100644
--- a/clang/include/clang/Basic/OpenMPKinds.h
+++ b/clang/include/clang/Basic/OpenMPKinds.h
@@ -217,6 +217,16 @@ enum OpenMPGrainsizeClauseModifier {
OMPC_GRAINSIZE_unknown
};
+enum OpenMPDynGroupprivateClauseModifier {
+#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name) OMPC_DYN_GROUPPRIVATE_##Name,
+#include "clang/Basic/OpenMPKinds.def"
+ 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"
diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h
index 91c3d4bd5210e..6e49c3a66d393 100644
--- a/clang/include/clang/Sema/SemaOpenMP.h
+++ b/clang/include/clang/Sema/SemaOpenMP.h
@@ -1385,6 +1385,12 @@ class SemaOpenMP : public SemaBase {
SourceLocation LParenLoc,
SourceLocation EndLoc);
+ /// Called on a well-formed 'dyn_groupprivate' clause.
+ 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 *
ActOnOpenMPDoacrossClause(OpenMPDoacrossClauseModifier DepType,
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index de8b5996818de..4ecf2355aac16 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -104,6 +104,8 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
return static_cast<const OMPFilterClause *>(C);
case OMPC_ompx_dyn_cgroup_mem:
return static_cast<const OMPXDynCGroupMemClause *>(C);
+ case OMPC_dyn_groupprivate:
+ return static_cast<const OMPDynGroupprivateClause *>(C);
case OMPC_default:
case OMPC_proc_bind:
case OMPC_safelen:
@@ -2725,6 +2727,22 @@ void OMPClausePrinter::VisitOMPXDynCGroupMemClause(
OS << ")";
}
+void OMPClausePrinter::VisitOMPDynGroupprivateClause(OMPDynGroupprivateClause *Node) {
+ OS << "dyn_groupprivate(";
+ 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 << ")";
+}
+
void OMPClausePrinter::VisitOMPDoacrossClause(OMPDoacrossClause *Node) {
OS << "doacross(";
OpenMPDoacrossClauseModifier DepType = Node->getDependenceType();
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index c61450e19f1b6..6b1a016649547 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -957,6 +957,12 @@ void OMPClauseProfiler::VisitOMPXDynCGroupMemClause(
if (Expr *Size = C->getSize())
Profiler->VisitStmt(Size);
}
+void OMPClauseProfiler::VisitOMPDynGroupprivateClause(
+ const OMPDynGroupprivateClause *C) {
+ VistOMPClauseWithPreInit(C);
+ if (auto *Size = C->getSize())
+ Profiler->VisitStmt(Size);
+}
void OMPClauseProfiler::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
VisitOMPClauseList(C);
}
diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp
index d3d393bd09396..64ecaf5c9813e 100644
--- a/clang/lib/Basic/OpenMPKinds.cpp
+++ b/clang/lib/Basic/OpenMPKinds.cpp
@@ -171,6 +171,12 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
return OMPC_GRAINSIZE_unknown;
return Type;
}
+ case OMPC_dyn_groupprivate: {
+ return llvm::StringSwitch<unsigned>(Str)
+#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name) .Case(#Name, OMPC_DYN_GROUPPRIVATE_##Name)
+#include "clang/Basic/OpenMPKinds.def"
+ .Default(OMPC_DYN_GROUPPRIVATE_unknown);
+ }
case OMPC_num_tasks: {
unsigned Type = llvm::StringSwitch<unsigned>(Str)
#define OPENMP_NUMTASKS_MODIFIER(Name) .Case(#Name, OMPC_NUMTASKS_##Name)
@@ -508,6 +514,16 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
#include "clang/Basic/OpenMPKinds.def"
}
llvm_unreachable("Invalid OpenMP 'grainsize' clause modifier");
+ case OMPC_dyn_groupprivate:
+ switch (Type) {
+ case OMPC_DYN_GROUPPRIVATE_unknown:
+ return "unknown";
+#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name) \
+ case OMPC_DYN_GROUPPRIVATE_##Name: \
+ return #Name;
+#include "clang/Basic/OpenMPKinds.def"
+ }
+ llvm_unreachable("Invalid OpenMP 'dyn_groupprivate' clause modifier");
case OMPC_num_tasks:
switch (Type) {
case OMPC_NUMTASKS_unknown:
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index a5f2f0efa2c3b..cb74e48a7c967 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9489,18 +9489,27 @@ static llvm::Value *emitDeviceID(
return DeviceID;
}
-static llvm::Value *emitDynCGGroupMem(const OMPExecutableDirective &D,
- CodeGenFunction &CGF) {
- llvm::Value *DynCGroupMem = CGF.Builder.getInt32(0);
-
- if (auto *DynMemClause = D.getSingleClause<OMPXDynCGroupMemClause>()) {
- CodeGenFunction::RunCleanupsScope DynCGroupMemScope(CGF);
- llvm::Value *DynCGroupMemVal = CGF.EmitScalarExpr(
- DynMemClause->getSize(), /*IgnoreResultAssign=*/true);
- DynCGroupMem = CGF.Builder.CreateIntCast(DynCGroupMemVal, CGF.Int32Ty,
- /*isSigned=*/false);
- }
- return DynCGroupMem;
+static std::pair<llvm::Value *, bool> emitDynCGroupMem(const OMPExecutableDirective &D,
+ CodeGenFunction &CGF) {
+ llvm::Value *DynGP = CGF.Builder.getInt32(0);
+ bool DynGPFallback = false;
+
+ if (auto *DynGPClause = D.getSingleClause<OMPDynGroupprivateClause>()) {
+ CodeGenFunction::RunCleanupsScope DynGPScope(CGF);
+ llvm::Value *DynGPVal = CGF.EmitScalarExpr(
+ DynGPClause->getSize(), /*IgnoreResultAssign=*/true);
+ DynGP = CGF.Builder.CreateIntCast(DynGPVal, CGF.Int32Ty,
+ /*isSigned=*/false);
+ 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(
+ OMPXDynCGClause->getSize(), /*IgnoreResultAssign=*/true);
+ DynGP = CGF.Builder.CreateIntCast(DynCGMemVal, CGF.Int32Ty,
+ /*isSigned=*/false);
+ }
+ return { DynGP, DynGPFallback };
}
static void genMapInfoForCaptures(
MappableExprsHandler &MEHandler, CodeGenFunction &CGF,
@@ -9710,7 +9719,7 @@ static void emitTargetCallKernelLaunch(
llvm::Value *RTLoc = OMPRuntime->emitUpdateLocation(CGF, D.getBeginLoc());
llvm::Value *NumIterations =
OMPRuntime->emitTargetNumIterationsCall(CGF, D, SizeEmitter);
- llvm::Value *DynCGGroupMem = emitDynCGGroupMem(D, CGF);
+ auto [DynCGroupMem, DynCGroupMemFallback] = emitDynCGroupMem(D, CGF);
llvm::OpenMPIRBuilder::InsertPointTy AllocaIP(
CGF.AllocaInsertPt->getParent(), CGF.AllocaInsertPt->getIterator());
@@ -9720,7 +9729,7 @@ static void emitTargetCallKernelLaunch(
llvm::OpenMPIRBuilder::TargetKernelArgs Args(
NumTargetItems, RTArgs, NumIterations, NumTeams, NumThreads,
- DynCGGroupMem, HasNoWait);
+ DynCGroupMem, HasNoWait, DynCGroupMemFallback);
llvm::OpenMPIRBuilder::InsertPointTy AfterIP =
cantFail(OMPRuntime->getOMPBuilder().emitKernelLaunch(
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index aa6a0c61a2c17..91eed315a6786 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -3039,6 +3039,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
case OMPC_align:
case OMPC_message:
case OMPC_ompx_dyn_cgroup_mem:
+ case OMPC_dyn_groupprivate:
// OpenMP [2.5, Restrictions]
// At most one num_threads clause can appear on the directive.
// OpenMP [2.8.1, simd construct, Restrictions]
@@ -3077,7 +3078,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
PP.LookAhead(/*N=*/0).isNot(tok::l_paren))
Clause = ParseOpenMPClause(CKind, WrongDirective);
else if (CKind == OMPC_grainsize || CKind == OMPC_num_tasks ||
- CKind == OMPC_num_threads)
+ CKind == OMPC_num_threads || CKind == OMPC_dyn_groupprivate)
Clause = ParseOpenMPSingleExprWithArgClause(DKind, CKind, WrongDirective);
else
Clause = ParseOpenMPSingleExprClause(CKind, WrongDirective);
@@ -3835,6 +3836,39 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
Arg.push_back(OMPC_GRAINSIZE_unknown);
KLoc.emplace_back();
}
+ } else if (Kind == OMPC_dyn_groupprivate) {
+ 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();
+ if (Tok.is(tok::comma)) {
+ // Parse ',' 'modifier'
+ ConsumeAnyToken();
+ 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();
+ }
+ // 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> ':'
OpenMPNumTasksClauseModifier Modifier =
@@ -3913,7 +3947,8 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
(Kind == OMPC_dist_schedule && DelimLoc.isValid()) ||
Kind == OMPC_if || Kind == OMPC_device ||
Kind == OMPC_grainsize || Kind == OMPC_num_tasks ||
- Kind == OMPC_num_threads;
+ Kind == OMPC_num_threads ||
+ Kind == OMPC_dyn_groupprivate;
if (NeedAnExpression) {
SourceLocation ELoc = Tok.getLocation();
ExprResult LHS(
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 4ecc9b0d4c5c8..d91c1adf1e18d 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@...
[truncated]
|
|
@llvm/pr-subscribers-clang-modules Author: Kevin Sala Penades (kevinsala) ChangesThis PR adds support for the The implementation reuses part of the runtime support for Patch is 92.16 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/152651.diff 45 Files Affected:
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 1118d3e062e68..2d6931ee052d8 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -9768,6 +9768,161 @@ class OMPXDynCGroupMemClause
Expr *getSize() const { return getStmtAs<Expr>(); }
};
+/// This represents 'dyn_groupprivate' clause in '#pragma omp target ...'
+/// and '#pragma omp teams ...' directives.
+///
+/// \code
+/// #pragma omp target [...] dyn_groupprivate(a,b: N)
+/// \endcode
+class OMPDynGroupprivateClause : public OMPClause, public OMPClauseWithPreInit {
+ friend class OMPClauseReader;
+
+ /// Location of '('.
+ SourceLocation LParenLoc;
+
+ /// Modifiers for 'dyn_groupprivate' clause.
+ enum {FIRST, SECOND, NUM_MODIFIERS};
+ OpenMPDynGroupprivateClauseModifier Modifiers[NUM_MODIFIERS];
+
+ /// Locations of modifiers.
+ SourceLocation ModifiersLoc[NUM_MODIFIERS];
+
+ /// The size of the dyn_groupprivate.
+ Expr *Size = nullptr;
+
+ /// 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 with a size expression \a Size.
+ ///
+ /// \param StartLoc Starting location of the clause.
+ /// \param LParenLoc Location of '('.
+ /// \param EndLoc Ending location of the clause.
+ /// \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.
+ explicit OMPDynGroupprivateClause()
+ : OMPClause(llvm::omp::OMPC_dyn_groupprivate, SourceLocation(), SourceLocation()),
+ OMPClauseWithPreInit(this) {
+ Modifiers[FIRST] = OMPC_DYN_GROUPPRIVATE_unknown;
+ Modifiers[SECOND] = OMPC_DYN_GROUPPRIVATE_unknown;
+ }
+
+ /// Get the first modifier of the clause.
+ OpenMPDynGroupprivateClauseModifier getFirstDynGroupprivateModifier() const {
+ return Modifiers[FIRST];
+ }
+
+ /// Get the second modifier of the clause.
+ OpenMPDynGroupprivateClauseModifier getSecondDynGroupprivateModifier() const {
+ return Modifiers[SECOND];
+ }
+
+ /// Get location of '('.
+ SourceLocation getLParenLoc() { return LParenLoc; }
+
+ /// 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'
/// directive.
///
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 5cb2f57edffe4..129115d56fe82 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -4060,6 +4060,14 @@ bool RecursiveASTVisitor<Derived>::VisitOMPXDynCGroupMemClause(
return true;
}
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPDynGroupprivateClause(
+ OMPDynGroupprivateClause *C) {
+ TRY_TO(VisitOMPClauseWithPreInit(C));
+ TRY_TO(TraverseStmt(C->getSize()));
+ return true;
+}
+
template <typename Derived>
bool RecursiveASTVisitor<Derived>::VisitOMPDoacrossClause(
OMPDoacrossClause *C) {
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index b285309e0b3ca..45b2a79d11561 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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<
diff --git a/clang/include/clang/Basic/OpenMPKinds.def b/clang/include/clang/Basic/OpenMPKinds.def
index 9d6f816eea91f..3321e19cae9b1 100644
--- a/clang/include/clang/Basic/OpenMPKinds.def
+++ b/clang/include/clang/Basic/OpenMPKinds.def
@@ -83,6 +83,9 @@
#ifndef OPENMP_GRAINSIZE_MODIFIER
#define OPENMP_GRAINSIZE_MODIFIER(Name)
#endif
+#ifndef OPENMP_DYN_GROUPPRIVATE_MODIFIER
+#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name)
+#endif
#ifndef OPENMP_NUMTASKS_MODIFIER
#define OPENMP_NUMTASKS_MODIFIER(Name)
#endif
@@ -227,6 +230,11 @@ OPENMP_BIND_KIND(thread)
// Modifiers for the 'grainsize' clause.
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)
+
// Modifiers for the 'num_tasks' clause.
OPENMP_NUMTASKS_MODIFIER(strict)
@@ -245,6 +253,7 @@ OPENMP_DOACROSS_MODIFIER(source_omp_cur_iteration)
#undef OPENMP_NUMTASKS_MODIFIER
#undef OPENMP_NUMTHREADS_MODIFIER
+#undef OPENMP_DYN_GROUPPRIVATE_MODIFIER
#undef OPENMP_GRAINSIZE_MODIFIER
#undef OPENMP_BIND_KIND
#undef OPENMP_ADJUST_ARGS_KIND
diff --git a/clang/include/clang/Basic/OpenMPKinds.h b/clang/include/clang/Basic/OpenMPKinds.h
index f40db4c13c55a..3e164bf1adf22 100644
--- a/clang/include/clang/Basic/OpenMPKinds.h
+++ b/clang/include/clang/Basic/OpenMPKinds.h
@@ -217,6 +217,16 @@ enum OpenMPGrainsizeClauseModifier {
OMPC_GRAINSIZE_unknown
};
+enum OpenMPDynGroupprivateClauseModifier {
+#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name) OMPC_DYN_GROUPPRIVATE_##Name,
+#include "clang/Basic/OpenMPKinds.def"
+ 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"
diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h
index 91c3d4bd5210e..6e49c3a66d393 100644
--- a/clang/include/clang/Sema/SemaOpenMP.h
+++ b/clang/include/clang/Sema/SemaOpenMP.h
@@ -1385,6 +1385,12 @@ class SemaOpenMP : public SemaBase {
SourceLocation LParenLoc,
SourceLocation EndLoc);
+ /// Called on a well-formed 'dyn_groupprivate' clause.
+ 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 *
ActOnOpenMPDoacrossClause(OpenMPDoacrossClauseModifier DepType,
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index de8b5996818de..4ecf2355aac16 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -104,6 +104,8 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
return static_cast<const OMPFilterClause *>(C);
case OMPC_ompx_dyn_cgroup_mem:
return static_cast<const OMPXDynCGroupMemClause *>(C);
+ case OMPC_dyn_groupprivate:
+ return static_cast<const OMPDynGroupprivateClause *>(C);
case OMPC_default:
case OMPC_proc_bind:
case OMPC_safelen:
@@ -2725,6 +2727,22 @@ void OMPClausePrinter::VisitOMPXDynCGroupMemClause(
OS << ")";
}
+void OMPClausePrinter::VisitOMPDynGroupprivateClause(OMPDynGroupprivateClause *Node) {
+ OS << "dyn_groupprivate(";
+ 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 << ")";
+}
+
void OMPClausePrinter::VisitOMPDoacrossClause(OMPDoacrossClause *Node) {
OS << "doacross(";
OpenMPDoacrossClauseModifier DepType = Node->getDependenceType();
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index c61450e19f1b6..6b1a016649547 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -957,6 +957,12 @@ void OMPClauseProfiler::VisitOMPXDynCGroupMemClause(
if (Expr *Size = C->getSize())
Profiler->VisitStmt(Size);
}
+void OMPClauseProfiler::VisitOMPDynGroupprivateClause(
+ const OMPDynGroupprivateClause *C) {
+ VistOMPClauseWithPreInit(C);
+ if (auto *Size = C->getSize())
+ Profiler->VisitStmt(Size);
+}
void OMPClauseProfiler::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
VisitOMPClauseList(C);
}
diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp
index d3d393bd09396..64ecaf5c9813e 100644
--- a/clang/lib/Basic/OpenMPKinds.cpp
+++ b/clang/lib/Basic/OpenMPKinds.cpp
@@ -171,6 +171,12 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
return OMPC_GRAINSIZE_unknown;
return Type;
}
+ case OMPC_dyn_groupprivate: {
+ return llvm::StringSwitch<unsigned>(Str)
+#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name) .Case(#Name, OMPC_DYN_GROUPPRIVATE_##Name)
+#include "clang/Basic/OpenMPKinds.def"
+ .Default(OMPC_DYN_GROUPPRIVATE_unknown);
+ }
case OMPC_num_tasks: {
unsigned Type = llvm::StringSwitch<unsigned>(Str)
#define OPENMP_NUMTASKS_MODIFIER(Name) .Case(#Name, OMPC_NUMTASKS_##Name)
@@ -508,6 +514,16 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
#include "clang/Basic/OpenMPKinds.def"
}
llvm_unreachable("Invalid OpenMP 'grainsize' clause modifier");
+ case OMPC_dyn_groupprivate:
+ switch (Type) {
+ case OMPC_DYN_GROUPPRIVATE_unknown:
+ return "unknown";
+#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name) \
+ case OMPC_DYN_GROUPPRIVATE_##Name: \
+ return #Name;
+#include "clang/Basic/OpenMPKinds.def"
+ }
+ llvm_unreachable("Invalid OpenMP 'dyn_groupprivate' clause modifier");
case OMPC_num_tasks:
switch (Type) {
case OMPC_NUMTASKS_unknown:
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index a5f2f0efa2c3b..cb74e48a7c967 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9489,18 +9489,27 @@ static llvm::Value *emitDeviceID(
return DeviceID;
}
-static llvm::Value *emitDynCGGroupMem(const OMPExecutableDirective &D,
- CodeGenFunction &CGF) {
- llvm::Value *DynCGroupMem = CGF.Builder.getInt32(0);
-
- if (auto *DynMemClause = D.getSingleClause<OMPXDynCGroupMemClause>()) {
- CodeGenFunction::RunCleanupsScope DynCGroupMemScope(CGF);
- llvm::Value *DynCGroupMemVal = CGF.EmitScalarExpr(
- DynMemClause->getSize(), /*IgnoreResultAssign=*/true);
- DynCGroupMem = CGF.Builder.CreateIntCast(DynCGroupMemVal, CGF.Int32Ty,
- /*isSigned=*/false);
- }
- return DynCGroupMem;
+static std::pair<llvm::Value *, bool> emitDynCGroupMem(const OMPExecutableDirective &D,
+ CodeGenFunction &CGF) {
+ llvm::Value *DynGP = CGF.Builder.getInt32(0);
+ bool DynGPFallback = false;
+
+ if (auto *DynGPClause = D.getSingleClause<OMPDynGroupprivateClause>()) {
+ CodeGenFunction::RunCleanupsScope DynGPScope(CGF);
+ llvm::Value *DynGPVal = CGF.EmitScalarExpr(
+ DynGPClause->getSize(), /*IgnoreResultAssign=*/true);
+ DynGP = CGF.Builder.CreateIntCast(DynGPVal, CGF.Int32Ty,
+ /*isSigned=*/false);
+ 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(
+ OMPXDynCGClause->getSize(), /*IgnoreResultAssign=*/true);
+ DynGP = CGF.Builder.CreateIntCast(DynCGMemVal, CGF.Int32Ty,
+ /*isSigned=*/false);
+ }
+ return { DynGP, DynGPFallback };
}
static void genMapInfoForCaptures(
MappableExprsHandler &MEHandler, CodeGenFunction &CGF,
@@ -9710,7 +9719,7 @@ static void emitTargetCallKernelLaunch(
llvm::Value *RTLoc = OMPRuntime->emitUpdateLocation(CGF, D.getBeginLoc());
llvm::Value *NumIterations =
OMPRuntime->emitTargetNumIterationsCall(CGF, D, SizeEmitter);
- llvm::Value *DynCGGroupMem = emitDynCGGroupMem(D, CGF);
+ auto [DynCGroupMem, DynCGroupMemFallback] = emitDynCGroupMem(D, CGF);
llvm::OpenMPIRBuilder::InsertPointTy AllocaIP(
CGF.AllocaInsertPt->getParent(), CGF.AllocaInsertPt->getIterator());
@@ -9720,7 +9729,7 @@ static void emitTargetCallKernelLaunch(
llvm::OpenMPIRBuilder::TargetKernelArgs Args(
NumTargetItems, RTArgs, NumIterations, NumTeams, NumThreads,
- DynCGGroupMem, HasNoWait);
+ DynCGroupMem, HasNoWait, DynCGroupMemFallback);
llvm::OpenMPIRBuilder::InsertPointTy AfterIP =
cantFail(OMPRuntime->getOMPBuilder().emitKernelLaunch(
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index aa6a0c61a2c17..91eed315a6786 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -3039,6 +3039,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
case OMPC_align:
case OMPC_message:
case OMPC_ompx_dyn_cgroup_mem:
+ case OMPC_dyn_groupprivate:
// OpenMP [2.5, Restrictions]
// At most one num_threads clause can appear on the directive.
// OpenMP [2.8.1, simd construct, Restrictions]
@@ -3077,7 +3078,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
PP.LookAhead(/*N=*/0).isNot(tok::l_paren))
Clause = ParseOpenMPClause(CKind, WrongDirective);
else if (CKind == OMPC_grainsize || CKind == OMPC_num_tasks ||
- CKind == OMPC_num_threads)
+ CKind == OMPC_num_threads || CKind == OMPC_dyn_groupprivate)
Clause = ParseOpenMPSingleExprWithArgClause(DKind, CKind, WrongDirective);
else
Clause = ParseOpenMPSingleExprClause(CKind, WrongDirective);
@@ -3835,6 +3836,39 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
Arg.push_back(OMPC_GRAINSIZE_unknown);
KLoc.emplace_back();
}
+ } else if (Kind == OMPC_dyn_groupprivate) {
+ 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();
+ if (Tok.is(tok::comma)) {
+ // Parse ',' 'modifier'
+ ConsumeAnyToken();
+ 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();
+ }
+ // 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> ':'
OpenMPNumTasksClauseModifier Modifier =
@@ -3913,7 +3947,8 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
(Kind == OMPC_dist_schedule && DelimLoc.isValid()) ||
Kind == OMPC_if || Kind == OMPC_device ||
Kind == OMPC_grainsize || Kind == OMPC_num_tasks ||
- Kind == OMPC_num_threads;
+ Kind == OMPC_num_threads ||
+ Kind == OMPC_dyn_groupprivate;
if (NeedAnExpression) {
SourceLocation ELoc = Tok.getLocation();
ExprResult LHS(
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 4ecc9b0d4c5c8..d91c1adf1e18d 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@...
[truncated]
|
arsenm
reviewed
Aug 8, 2025
mjklemm
reviewed
Aug 8, 2025
kevinsala
force-pushed
the
users/kevinsala/omp-dyn-groupprivate-pr
branch
from
184b0ae to
099c502
Compare
kevinsala
changed the title
This was referenced Aug 9, 2025
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR adds support for the
dyn_groupprivateclause, which will be part of OpenMP 6.1. This feature allows users to request dynamic shared memory on target regions. Runtime support for host constructs (e.g.,#pragma omp teams) is not implemented yet.The implementation reuses part of the runtime support for
ompx_dyn_cgroup_memclause.This PR is the part 1, which adds the parser/semantic support in Clang. Part 2 is #152830 (codegen support) and part 3 is #152831 (offload runtime support).