[clang] [llvm] [OpenMP][Offload] Add support for dyn_groupprivate clause (PR #152651)
Kevin Sala Penades via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 8 18:34:50 PDT 2025
https://github.com/kevinsala updated https://github.com/llvm/llvm-project/pull/152651
>From 099c502bdf02ed9bc34bbfc70a6e786746ecee90 Mon Sep 17 00:00:00 2001
From: Kevin Sala <salapenades1 at llnl.gov>
Date: Fri, 8 Aug 2025 10:43:52 -0700
Subject: [PATCH] [OpenMP] Add parser/semantic support for dyn_groupprivate
clause
---
clang/include/clang/AST/OpenMPClause.h | 155 ++++++++++++++++++
clang/include/clang/AST/RecursiveASTVisitor.h | 8 +
.../clang/Basic/DiagnosticSemaKinds.td | 3 +
clang/include/clang/Basic/OpenMPKinds.def | 9 +
clang/include/clang/Basic/OpenMPKinds.h | 10 ++
clang/include/clang/Sema/SemaOpenMP.h | 7 +
clang/lib/AST/OpenMPClause.cpp | 21 +++
clang/lib/AST/StmtProfile.cpp | 6 +
clang/lib/Basic/OpenMPKinds.cpp | 17 ++
clang/lib/Parse/ParseOpenMP.cpp | 47 +++++-
clang/lib/Sema/SemaOpenMP.cpp | 73 ++++++++-
clang/lib/Sema/TreeTransform.h | 26 +++
clang/lib/Serialization/ASTReader.cpp | 16 ++
clang/lib/Serialization/ASTWriter.cpp | 11 ++
.../target_dyn_groupprivate_messages.cpp | 87 ++++++++++
...target_teams_dyn_groupprivate_messages.cpp | 87 ++++++++++
clang/tools/libclang/CIndex.cpp | 5 +
llvm/include/llvm/Frontend/OpenMP/OMP.td | 25 +++
18 files changed, 605 insertions(+), 8 deletions(-)
create mode 100644 clang/test/OpenMP/target_dyn_groupprivate_messages.cpp
create mode 100644 clang/test/OpenMP/target_teams_dyn_groupprivate_messages.cpp
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 1118d3e062e68..a3983120df069 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..edfa2d229789a 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11995,6 +11995,9 @@ 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..3b161ff3c7d45 100644
--- a/clang/include/clang/Sema/SemaOpenMP.h
+++ b/clang/include/clang/Sema/SemaOpenMP.h
@@ -1385,6 +1385,13 @@ 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..a6e79f1e2230b 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,25 @@ void OMPClausePrinter::VisitOMPXDynCGroupMemClause(
OS << ")";
}
+void OMPClausePrinter::VisitOMPDynGroupprivateClause(
+ OMPDynGroupprivateClause *Node) {
+ OS << "dyn_groupprivate(";
+ if (Node->getFirstDynGroupprivateModifier() !=
+ OMPC_DYN_GROUPPRIVATE_unknown) {
+ OS << getOpenMPSimpleClauseTypeName(
+ OMPC_dyn_groupprivate, Node->getFirstDynGroupprivateModifier());
+ if (Node->getSecondDynGroupprivateModifier() !=
+ OMPC_DYN_GROUPPRIVATE_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..5a2b578cfb33e 100644
--- a/clang/lib/Basic/OpenMPKinds.cpp
+++ b/clang/lib/Basic/OpenMPKinds.cpp
@@ -171,6 +171,13 @@ 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 +515,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/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index aa6a0c61a2c17..fe3c20b765e52 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,40 @@ 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 =
@@ -3909,11 +3944,11 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
}
}
- bool NeedAnExpression = (Kind == OMPC_schedule && DelimLoc.isValid()) ||
- (Kind == OMPC_dist_schedule && DelimLoc.isValid()) ||
- Kind == OMPC_if || Kind == OMPC_device ||
- Kind == OMPC_grainsize || Kind == OMPC_num_tasks ||
- Kind == OMPC_num_threads;
+ bool NeedAnExpression =
+ (Kind == OMPC_schedule && DelimLoc.isValid()) ||
+ (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_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..c63310757513e 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -15560,6 +15560,7 @@ OMPClause *SemaOpenMP::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind,
case OMPC_holds:
Res = ActOnOpenMPHoldsClause(Expr, StartLoc, LParenLoc, EndLoc);
break;
+ case OMPC_dyn_groupprivate:
case OMPC_grainsize:
case OMPC_num_tasks:
case OMPC_num_threads:
@@ -15686,6 +15687,8 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
case OMPC_num_teams:
case OMPC_thread_limit:
case OMPC_ompx_dyn_cgroup_mem:
+ case OMPC_dyn_groupprivate:
+ // TODO: This may need to consider teams too.
if (Leafs[0] == OMPD_target)
return OMPD_target;
break;
@@ -16646,7 +16649,7 @@ OMPClause *SemaOpenMP::ActOnOpenMPSingleExprWithArgClause(
SourceLocation EndLoc) {
OMPClause *Res = nullptr;
switch (Kind) {
- case OMPC_schedule:
+ case OMPC_schedule: {
enum { Modifier1, Modifier2, ScheduleKind, NumberOfElements };
assert(Argument.size() == NumberOfElements &&
ArgumentLoc.size() == NumberOfElements);
@@ -16656,7 +16659,7 @@ OMPClause *SemaOpenMP::ActOnOpenMPSingleExprWithArgClause(
static_cast<OpenMPScheduleClauseKind>(Argument[ScheduleKind]), Expr,
StartLoc, LParenLoc, ArgumentLoc[Modifier1], ArgumentLoc[Modifier2],
ArgumentLoc[ScheduleKind], DelimLoc, EndLoc);
- break;
+ } break;
case OMPC_if:
assert(Argument.size() == 1 && ArgumentLoc.size() == 1);
Res = ActOnOpenMPIfClause(static_cast<OpenMPDirectiveKind>(Argument.back()),
@@ -16703,6 +16706,16 @@ OMPClause *SemaOpenMP::ActOnOpenMPSingleExprWithArgClause(
static_cast<OpenMPNumTasksClauseModifier>(Argument.back()), Expr,
StartLoc, LParenLoc, ArgumentLoc.back(), EndLoc);
break;
+ case OMPC_dyn_groupprivate: {
+ enum { Modifier1, Modifier2, NumberOfElements };
+ assert(Argument.size() == NumberOfElements &&
+ ArgumentLoc.size() == NumberOfElements);
+ Res = ActOnOpenMPDynGroupprivateClause(
+ static_cast<OpenMPDynGroupprivateClauseModifier>(Argument[Modifier1]),
+ static_cast<OpenMPDynGroupprivateClauseModifier>(Argument[Modifier2]),
+ Expr, StartLoc, LParenLoc, ArgumentLoc[Modifier1],
+ ArgumentLoc[Modifier2], EndLoc);
+ } break;
case OMPC_num_threads:
assert(Argument.size() == 1 && ArgumentLoc.size() == 1 &&
"Modifier for num_threads clause and its location are expected.");
@@ -17056,6 +17069,7 @@ OMPClause *SemaOpenMP::ActOnOpenMPClause(OpenMPClauseKind Kind,
case OMPC_affinity:
case OMPC_when:
case OMPC_ompx_dyn_cgroup_mem:
+ case OMPC_dyn_groupprivate:
default:
llvm_unreachable("Clause is not allowed.");
}
@@ -24143,6 +24157,61 @@ OMPClause *SemaOpenMP::ActOnOpenMPXDynCGroupMemClause(Expr *Size,
ValExpr, HelperValStmt, CaptureRegion, StartLoc, LParenLoc, EndLoc);
}
+OMPClause *SemaOpenMP::ActOnOpenMPDynGroupprivateClause(
+ OpenMPDynGroupprivateClauseModifier M1,
+ OpenMPDynGroupprivateClauseModifier M2, Expr *Size, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation M1Loc, SourceLocation M2Loc,
+ SourceLocation EndLoc) {
+
+ if ((M1Loc.isValid() && M1 == OMPC_DYN_GROUPPRIVATE_unknown) ||
+ (M2Loc.isValid() && M2 == OMPC_DYN_GROUPPRIVATE_unknown)) {
+ std::string Values = getListOfPossibleValues(
+ OMPC_dyn_groupprivate, /*First=*/0, OMPC_DYN_GROUPPRIVATE_unknown);
+ Diag((M1Loc.isValid() && M1 == OMPC_DYN_GROUPPRIVATE_unknown) ? M1Loc
+ : M2Loc,
+ diag::err_omp_unexpected_clause_value)
+ << Values << getOpenMPClauseName(OMPC_dyn_groupprivate);
+ return nullptr;
+ }
+
+ if ((M1Loc.isValid() && M2Loc.isValid() && M1 == M2) ||
+ (M1 == OMPC_DYN_GROUPPRIVATE_strict &&
+ M2 == OMPC_DYN_GROUPPRIVATE_fallback) ||
+ (M1 == OMPC_DYN_GROUPPRIVATE_fallback &&
+ M2 == OMPC_DYN_GROUPPRIVATE_strict)) {
+
+ Diag(M2Loc, diag::err_omp_unexpected_dyn_groupprivate_modifier)
+ << getOpenMPSimpleClauseTypeName(OMPC_dyn_groupprivate, M2)
+ << getOpenMPSimpleClauseTypeName(OMPC_dyn_groupprivate, M1);
+ return nullptr;
+ }
+
+ Expr *ValExpr = Size;
+ Stmt *HelperValStmt = nullptr;
+
+ // OpenMP [2.5, Restrictions]
+ // The dyn_groupprivate expression must evaluate to a positive integer
+ // value.
+ if (!isNonNegativeIntegerValue(ValExpr, SemaRef, OMPC_dyn_groupprivate,
+ /*StrictlyPositive=*/false))
+ return nullptr;
+
+ OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective();
+ OpenMPDirectiveKind CaptureRegion = getOpenMPCaptureRegionForClause(
+ DKind, OMPC_dyn_groupprivate, getLangOpts().OpenMP);
+ if (CaptureRegion != OMPD_unknown &&
+ !SemaRef.CurContext->isDependentContext()) {
+ ValExpr = SemaRef.MakeFullExpr(ValExpr).get();
+ llvm::MapVector<const Expr *, DeclRefExpr *> Captures;
+ ValExpr = tryBuildCapture(SemaRef, ValExpr, Captures).get();
+ HelperValStmt = buildPreInits(getASTContext(), Captures);
+ }
+
+ return new (getASTContext()) OMPDynGroupprivateClause(
+ StartLoc, LParenLoc, EndLoc, ValExpr, HelperValStmt, CaptureRegion, M1,
+ M1Loc, M2, M2Loc);
+}
+
OMPClause *SemaOpenMP::ActOnOpenMPDoacrossClause(
OpenMPDoacrossClauseModifier DepType, SourceLocation DepLoc,
SourceLocation ColonLoc, ArrayRef<Expr *> VarList, SourceLocation StartLoc,
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 758012f894a41..6222d63374824 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -2430,6 +2430,19 @@ class TreeTransform {
LParenLoc, EndLoc);
}
+ /// Build a new OpenMP 'dyn_groupprivate' clause.
+ ///
+ /// By default, performs semantic analysis to build the new OpenMP clause.
+ /// Subclasses may override this routine to provide different behavior.
+ OMPClause *RebuildOMPDynGroupprivateClause(
+ OpenMPDynGroupprivateClauseModifier M1,
+ OpenMPDynGroupprivateClauseModifier M2, Expr *Size,
+ SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation M1Loc,
+ SourceLocation M2Loc, SourceLocation EndLoc) {
+ return getSema().OpenMP().ActOnOpenMPDynGroupprivateClause(
+ M1, M2, Size, StartLoc, LParenLoc, M1Loc, M2Loc, EndLoc);
+ }
+
/// Build a new OpenMP 'ompx_attribute' clause.
///
/// By default, performs semantic analysis to build the new OpenMP clause.
@@ -11691,6 +11704,19 @@ OMPClause *TreeTransform<Derived>::TransformOMPXDynCGroupMemClause(
Size.get(), C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
}
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPDynGroupprivateClause(
+ OMPDynGroupprivateClause *C) {
+ ExprResult Size = getDerived().TransformExpr(C->getSize());
+ if (Size.isInvalid())
+ return nullptr;
+ return getDerived().RebuildOMPDynGroupprivateClause(
+ C->getFirstDynGroupprivateModifier(),
+ C->getSecondDynGroupprivateModifier(), Size.get(), C->getBeginLoc(),
+ C->getLParenLoc(), C->getFirstDynGroupprivateModifierLoc(),
+ C->getSecondDynGroupprivateModifierLoc(), C->getEndLoc());
+}
+
template <typename Derived>
OMPClause *
TreeTransform<Derived>::TransformOMPDoacrossClause(OMPDoacrossClause *C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 30e0973149594..84a30fcfcea59 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11417,6 +11417,9 @@ OMPClause *OMPClauseReader::readClause() {
case llvm::omp::OMPC_ompx_dyn_cgroup_mem:
C = new (Context) OMPXDynCGroupMemClause();
break;
+ case llvm::omp::OMPC_dyn_groupprivate:
+ C = new (Context) OMPDynGroupprivateClause();
+ break;
case llvm::omp::OMPC_doacross: {
unsigned NumVars = Record.readInt();
unsigned NumLoops = Record.readInt();
@@ -12583,6 +12586,19 @@ void OMPClauseReader::VisitOMPXDynCGroupMemClause(OMPXDynCGroupMemClause *C) {
C->setLParenLoc(Record.readSourceLocation());
}
+void OMPClauseReader::VisitOMPDynGroupprivateClause(
+ OMPDynGroupprivateClause *C) {
+ VisitOMPClauseWithPreInit(C);
+ C->setFirstDynGroupprivateModifier(
+ static_cast<OpenMPDynGroupprivateClauseModifier>(Record.readInt()));
+ C->setSecondDynGroupprivateModifier(
+ static_cast<OpenMPDynGroupprivateClauseModifier>(Record.readInt()));
+ C->setSize(Record.readSubExpr());
+ C->setLParenLoc(Record.readSourceLocation());
+ C->setFirstDynGroupprivateModifierLoc(Record.readSourceLocation());
+ C->setSecondDynGroupprivateModifierLoc(Record.readSourceLocation());
+}
+
void OMPClauseReader::VisitOMPDoacrossClause(OMPDoacrossClause *C) {
C->setLParenLoc(Record.readSourceLocation());
C->setDependenceType(
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 1a20fc9595dce..42d1ab91a6879 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8586,6 +8586,17 @@ void OMPClauseWriter::VisitOMPXDynCGroupMemClause(OMPXDynCGroupMemClause *C) {
Record.AddSourceLocation(C->getLParenLoc());
}
+void OMPClauseWriter::VisitOMPDynGroupprivateClause(
+ OMPDynGroupprivateClause *C) {
+ VisitOMPClauseWithPreInit(C);
+ Record.push_back(C->getFirstDynGroupprivateModifier());
+ Record.push_back(C->getSecondDynGroupprivateModifier());
+ Record.AddStmt(C->getSize());
+ Record.AddSourceLocation(C->getLParenLoc());
+ Record.AddSourceLocation(C->getFirstDynGroupprivateModifierLoc());
+ Record.AddSourceLocation(C->getSecondDynGroupprivateModifierLoc());
+}
+
void OMPClauseWriter::VisitOMPDoacrossClause(OMPDoacrossClause *C) {
Record.push_back(C->varlist_size());
Record.push_back(C->getNumLoops());
diff --git a/clang/test/OpenMP/target_dyn_groupprivate_messages.cpp b/clang/test/OpenMP/target_dyn_groupprivate_messages.cpp
new file mode 100644
index 0000000000000..d5d855ee33e1f
--- /dev/null
+++ b/clang/test/OpenMP/target_dyn_groupprivate_messages.cpp
@@ -0,0 +1,87 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+template <class T, class S> // expected-note {{declared here}}
+int tmain(T argc, S **argv) {
+ T z;
+ #pragma omp target dyn_groupprivate // expected-error {{expected '(' after 'dyn_groupprivate'}}
+ foo();
+ #pragma omp target dyn_groupprivate ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate () // expected-error {{expected expression}}
+ foo();
+ #pragma omp target dyn_groupprivate (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
+ foo();
+ #pragma omp target dyn_groupprivate (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+ foo();
+ #pragma omp target dyn_groupprivate (foobool(argc)), dyn_groupprivate (true) // expected-error {{directive '#pragma omp target' cannot contain more than one 'dyn_groupprivate' clause}}
+ foo();
+ #pragma omp target dyn_groupprivate (S) // expected-error {{'S' does not refer to a value}}
+ foo();
+ #pragma omp target dyn_groupprivate (argv[1]=2) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate(argc+z)
+ foo();
+ return 0;
+}
+
+int main(int argc, char **argv) {
+constexpr int n = -1;
+int z;
+ #pragma omp target dyn_groupprivate // expected-error {{expected '(' after 'dyn_groupprivate'}}
+ foo();
+ #pragma omp target dyn_groupprivate ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate () // expected-error {{expected expression}}
+ foo();
+ #pragma omp target dyn_groupprivate (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
+ foo();
+ #pragma omp target dyn_groupprivate (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+ foo();
+ #pragma omp target dyn_groupprivate (foobool(argc)), dyn_groupprivate (true) // expected-error {{directive '#pragma omp target' cannot contain more than one 'dyn_groupprivate' clause}}
+ foo();
+ #pragma omp target dyn_groupprivate (S1) // expected-error {{'S1' does not refer to a value}}
+ foo();
+ #pragma omp target dyn_groupprivate (argv[1]=2) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate (1 0) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate(dyn_groupprivate(tmain(argc, argv) // expected-error2 {{expected ')'}} expected-note2 {{to match this '('}} expected-note {{in instantiation of function template specialization 'tmain<int, char>' requested here}}
+ foo();
+ #pragma omp target dyn_groupprivate(-1) // expected-error {{argument to 'dyn_groupprivate' clause must be a non-negative integer value}}
+ foo();
+ #pragma omp target dyn_groupprivate(cgrou) // expected-error {{use of undeclared identifier 'cgrou'}}
+ foo();
+ #pragma omp target dyn_groupprivate(cgrou: argc) // expected-error {{use of undeclared identifier 'cgrou'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target dyn_groupprivate(cgroup,cgroup: argc) // expected-error {{modifier 'cgroup' cannot be used along with modifier 'cgroup' in dyn_groupprivate}}
+ foo();
+ #pragma omp target dyn_groupprivate(fallback,strict: argc) // expected-error {{modifier 'strict' cannot be used along with modifier 'fallback' in dyn_groupprivate}}
+ foo();
+ #pragma omp target dyn_groupprivate(strict,fallback: argc) // expected-error {{modifier 'fallback' cannot be used along with modifier 'strict' in dyn_groupprivate}}
+ foo();
+ #pragma omp target dyn_groupprivate(: argc) // expected-error {{expected ')'}} expected-error {{expected expression}} expected-note {{to match this '('}}
+ foo();
+
+ return tmain(argc, argv);
+}
+
diff --git a/clang/test/OpenMP/target_teams_dyn_groupprivate_messages.cpp b/clang/test/OpenMP/target_teams_dyn_groupprivate_messages.cpp
new file mode 100644
index 0000000000000..422dff547355c
--- /dev/null
+++ b/clang/test/OpenMP/target_teams_dyn_groupprivate_messages.cpp
@@ -0,0 +1,87 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+
+void foo() {
+}
+
+bool foobool(int argc) {
+ return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+template <class T, class S> // expected-note {{declared here}}
+int tmain(T argc, S **argv) {
+ T z;
+ #pragma omp target teams dyn_groupprivate // expected-error {{expected '(' after 'dyn_groupprivate'}}
+ foo();
+ #pragma omp target teams dyn_groupprivate ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate () // expected-error {{expected expression}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (foobool(argc)), dyn_groupprivate (true) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'dyn_groupprivate' clause}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (S) // expected-error {{'S' does not refer to a value}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argv[1]=2) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate(argc+z)
+ foo();
+ return 0;
+}
+
+int main(int argc, char **argv) {
+constexpr int n = -1;
+int z;
+ #pragma omp target teams dyn_groupprivate // expected-error {{expected '(' after 'dyn_groupprivate'}}
+ foo();
+ #pragma omp target teams dyn_groupprivate ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate () // expected-error {{expected expression}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (foobool(argc)), dyn_groupprivate (true) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'dyn_groupprivate' clause}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (S1) // expected-error {{'S1' does not refer to a value}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argv[1]=2) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate (1 0) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate(dyn_groupprivate(tmain(argc, argv) // expected-error2 {{expected ')'}} expected-note2 {{to match this '('}} expected-note {{in instantiation of function template specialization 'tmain<int, char>' requested here}}
+ foo();
+ #pragma omp target teams dyn_groupprivate(-1) // expected-error {{argument to 'dyn_groupprivate' clause must be a non-negative integer value}}
+ foo();
+ #pragma omp target teams dyn_groupprivate(cgrou) // expected-error {{use of undeclared identifier 'cgrou'}}
+ foo();
+ #pragma omp target teams dyn_groupprivate(cgrou: argc) // expected-error {{use of undeclared identifier 'cgrou'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+ #pragma omp target teams dyn_groupprivate(cgroup,cgroup: argc) // expected-error {{modifier 'cgroup' cannot be used along with modifier 'cgroup' in dyn_groupprivate}}
+ foo();
+ #pragma omp target teams dyn_groupprivate(fallback,strict: argc) // expected-error {{modifier 'strict' cannot be used along with modifier 'fallback' in dyn_groupprivate}}
+ foo();
+ #pragma omp target teams dyn_groupprivate(strict,fallback: argc) // expected-error {{modifier 'fallback' cannot be used along with modifier 'strict' in dyn_groupprivate}}
+ foo();
+ #pragma omp target teams dyn_groupprivate(: argc) // expected-error {{expected ')'}} expected-error {{expected expression}} expected-note {{to match this '('}}
+ foo();
+
+ return tmain(argc, argv);
+}
+
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 9089984fa4a54..32d75d4cd8d38 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2808,6 +2808,11 @@ void OMPClauseEnqueue::VisitOMPXDynCGroupMemClause(
VisitOMPClauseWithPreInit(C);
Visitor->AddStmt(C->getSize());
}
+void OMPClauseEnqueue::VisitOMPDynGroupprivateClause(
+ const OMPDynGroupprivateClause *C) {
+ VisitOMPClauseWithPreInit(C);
+ Visitor->AddStmt(C->getSize());
+}
void OMPClauseEnqueue::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
VisitOMPClauseList(C);
}
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td
index 1b94657dfae1e..78442a63b3ac4 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMP.td
+++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td
@@ -178,6 +178,9 @@ def OMPC_Doacross : Clause<[Spelling<"doacross">]> {
def OMPC_DynamicAllocators : Clause<[Spelling<"dynamic_allocators">]> {
let clangClass = "OMPDynamicAllocatorsClause";
}
+def OMPC_DynGroupprivate : Clause<[Spelling<"dyn_groupprivate">]> {
+ let clangClass = "OMPDynGroupprivateClause";
+}
def OMPC_Enter : Clause<[Spelling<"enter">]> {
let flangClass = "OmpObjectList";
}
@@ -1104,6 +1107,7 @@ def OMP_Target : Directive<[Spelling<"target">]> {
let allowedOnceClauses = [
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_If>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_OMPX_Bare>,
@@ -1254,6 +1258,7 @@ def OMP_Teams : Directive<[Spelling<"teams">]> {
];
let allowedOnceClauses = [
VersionedClause<OMPC_Default>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_If, 52>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_ThreadLimit>,
@@ -1522,6 +1527,7 @@ def OMP_target_loop : Directive<[Spelling<"target loop">]> {
let allowedOnceClauses = [
VersionedClause<OMPC_Bind, 50>,
VersionedClause<OMPC_Collapse>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_Order>,
VersionedClause<OMPC_ThreadLimit>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
@@ -1983,6 +1989,7 @@ def OMP_TargetParallel : Directive<[Spelling<"target parallel">]> {
let allowedOnceClauses = [
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_ProcBind>,
@@ -2086,6 +2093,7 @@ def OMP_TargetParallelFor : Directive<[Spelling<"target parallel for">]> {
VersionedClause<OMPC_UsesAllocators, 50>,
];
let allowedOnceClauses = [
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_ThreadLimit, 51>,
];
@@ -2126,6 +2134,7 @@ def OMP_TargetParallelForSimd
VersionedClause<OMPC_UsesAllocators, 50>,
];
let allowedOnceClauses = [
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_ThreadLimit, 51>,
];
@@ -2155,6 +2164,7 @@ def OMP_target_parallel_loop : Directive<[Spelling<"target parallel loop">]> {
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_Default>,
VersionedClause<OMPC_DefaultMap>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
@@ -2189,6 +2199,7 @@ def OMP_TargetSimd : Directive<[Spelling<"target simd">]> {
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_Order, 50>,
@@ -2220,6 +2231,7 @@ def OMP_TargetTeams : Directive<[Spelling<"target teams">]> {
VersionedClause<OMPC_Default>,
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
@@ -2252,6 +2264,7 @@ def OMP_TargetTeamsDistribute
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_DistSchedule>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
@@ -2284,6 +2297,7 @@ def OMP_TargetTeamsDistributeParallelDo
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_DistSchedule>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_NumThreads>,
@@ -2322,6 +2336,7 @@ def OMP_TargetTeamsDistributeParallelDoSimd
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_DistSchedule>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_NumThreads>,
@@ -2367,6 +2382,7 @@ def OMP_TargetTeamsDistributeParallelFor
VersionedClause<OMPC_UsesAllocators, 50>,
];
let allowedOnceClauses = [
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
];
let leafConstructs =
@@ -2409,6 +2425,7 @@ def OMP_TargetTeamsDistributeParallelForSimd
VersionedClause<OMPC_UsesAllocators, 50>,
];
let allowedOnceClauses = [
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
];
let leafConstructs =
@@ -2441,6 +2458,7 @@ def OMP_TargetTeamsDistributeSimd
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_DistSchedule>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
@@ -2474,6 +2492,7 @@ def OMP_target_teams_loop : Directive<[Spelling<"target teams loop">]> {
VersionedClause<OMPC_Bind, 50>,
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_Default>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
@@ -2532,6 +2551,7 @@ def OMP_TeamsDistribute : Directive<[Spelling<"teams distribute">]> {
VersionedClause<OMPC_ThreadLimit>,
];
let allowedOnceClauses = [
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_If>,
VersionedClause<OMPC_Order, 50>,
];
@@ -2555,6 +2575,7 @@ def OMP_TeamsDistributeParallelDo
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_Default>,
VersionedClause<OMPC_DistSchedule>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_Order, 50>,
@@ -2584,6 +2605,7 @@ def OMP_TeamsDistributeParallelDoSimd
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_Default>,
VersionedClause<OMPC_DistSchedule>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_Order, 50>,
@@ -2632,6 +2654,7 @@ def OMP_TeamsDistributeParallelForSimd
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_Default>,
VersionedClause<OMPC_DistSchedule>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_If>,
VersionedClause<OMPC_LastPrivate>,
@@ -2673,6 +2696,7 @@ def OMP_TeamsDistributeSimd : Directive<[Spelling<"teams distribute simd">]> {
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_Default>,
VersionedClause<OMPC_DistSchedule>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_Order, 50>,
VersionedClause<OMPC_SafeLen>,
@@ -2696,6 +2720,7 @@ def OMP_teams_loop : Directive<[Spelling<"teams loop">]> {
VersionedClause<OMPC_Bind, 50>,
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_Default>,
+ VersionedClause<OMPC_DynGroupprivate>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_Order>,
VersionedClause<OMPC_ThreadLimit>,
More information about the llvm-commits
mailing list