Skip to content

Commit eb44bba

Browse files
kevinsalakevinsala
committed
[OpenMP][Offload] Add support for dyn_groupprivate clause
1 parent bc2cc8b commit eb44bba

File tree

41 files changed

+601
-84
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

41 files changed

+601
-84
lines changed

‎clang/include/clang/AST/OpenMPClause.h

Lines changed: 63 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -9768,6 +9768,69 @@ class OMPXDynCGroupMemClause
97689768
Expr *getSize() const { return getStmtAs<Expr>(); }
97699769
};
97709770

9771+
/// This represents 'dyn_groupprivate' clause in '#pragma omp target ...'
9772+
/// and '#pragma omp teams ...' directives.
9773+
///
9774+
/// \code
9775+
/// #pragma omp target [...] dyn_groupprivate(N)
9776+
/// \endcode
9777+
class OMPDynGroupprivateClause
9778+
: public OMPOneStmtClause<llvm::omp::OMPC_dyn_groupprivate, OMPClause>,
9779+
public OMPClauseWithPreInit {
9780+
friend class OMPClauseReader;
9781+
9782+
/// Modifiers for 'grainsize' clause.
9783+
OpenMPDynGroupprivateClauseModifier Modifier = OMPC_DYN_GROUPPRIVATE_unknown;
9784+
9785+
/// Location of the modifier.
9786+
SourceLocation ModifierLoc;
9787+
9788+
/// Set size.
9789+
void setSize(Expr *E) { setStmt(E); }
9790+
9791+
/// Sets modifier.
9792+
void setModifier(OpenMPDynGroupprivateClauseModifier M) { Modifier = M; }
9793+
9794+
/// Sets modifier location.
9795+
void setModifierLoc(SourceLocation Loc) { ModifierLoc = Loc; }
9796+
9797+
public:
9798+
/// Build 'dyn_groupprivate' clause.
9799+
///
9800+
/// \param Modifier Clause modifier.
9801+
/// \param Size Size expression.
9802+
/// \param HelperSize Helper Size expression
9803+
/// \param CaptureRegion Innermost OpenMP region where expressions in this
9804+
/// \param StartLoc Starting location of the clause.
9805+
/// \param LParenLoc Location of '('.
9806+
/// \param ModifierLoc Modifier location.
9807+
/// \param EndLoc Ending location of the clause.
9808+
OMPDynGroupprivateClause(OpenMPDynGroupprivateClauseModifier Modifier,
9809+
Expr *Size, Stmt *HelperSize,
9810+
OpenMPDirectiveKind CaptureRegion,
9811+
SourceLocation StartLoc, SourceLocation LParenLoc,
9812+
SourceLocation ModifierLoc, SourceLocation EndLoc)
9813+
: OMPOneStmtClause(Size, StartLoc, LParenLoc, EndLoc),
9814+
OMPClauseWithPreInit(this), Modifier(Modifier), ModifierLoc(ModifierLoc) {
9815+
setPreInitStmt(HelperSize, CaptureRegion);
9816+
}
9817+
9818+
/// Build an empty clause.
9819+
OMPDynGroupprivateClause() : OMPOneStmtClause(), OMPClauseWithPreInit(this) {}
9820+
9821+
/// Return the size expression.
9822+
Expr *getSize() { return getStmtAs<Expr>(); }
9823+
9824+
/// Return the size expression.
9825+
Expr *getSize() const { return getStmtAs<Expr>(); }
9826+
9827+
/// Gets modifier.
9828+
OpenMPDynGroupprivateClauseModifier getModifier() const { return Modifier; }
9829+
9830+
/// Gets modifier location.
9831+
SourceLocation getModifierLoc() const { return ModifierLoc; }
9832+
};
9833+
97719834
/// This represents the 'doacross' clause for the '#pragma omp ordered'
97729835
/// directive.
97739836
///

‎clang/include/clang/AST/RecursiveASTVisitor.h

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4060,6 +4060,14 @@ bool RecursiveASTVisitor<Derived>::VisitOMPXDynCGroupMemClause(
40604060
return true;
40614061
}
40624062

4063+
template <typename Derived>
4064+
bool RecursiveASTVisitor<Derived>::VisitOMPDynGroupprivateClause(
4065+
OMPDynGroupprivateClause *C) {
4066+
TRY_TO(VisitOMPClauseWithPreInit(C));
4067+
TRY_TO(TraverseStmt(C->getSize()));
4068+
return true;
4069+
}
4070+
40634071
template <typename Derived>
40644072
bool RecursiveASTVisitor<Derived>::VisitOMPDoacrossClause(
40654073
OMPDoacrossClause *C) {

‎clang/include/clang/Basic/OpenMPKinds.def

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -83,6 +83,9 @@
8383
#ifndef OPENMP_GRAINSIZE_MODIFIER
8484
#define OPENMP_GRAINSIZE_MODIFIER(Name)
8585
#endif
86+
#ifndef OPENMP_DYN_GROUPPRIVATE_MODIFIER
87+
#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name)
88+
#endif
8689
#ifndef OPENMP_NUMTASKS_MODIFIER
8790
#define OPENMP_NUMTASKS_MODIFIER(Name)
8891
#endif
@@ -227,6 +230,10 @@ OPENMP_BIND_KIND(thread)
227230
// Modifiers for the 'grainsize' clause.
228231
OPENMP_GRAINSIZE_MODIFIER(strict)
229232

233+
// Modifiers for the 'dyn_groupprivate' clause.
234+
OPENMP_DYN_GROUPPRIVATE_MODIFIER(strict)
235+
OPENMP_DYN_GROUPPRIVATE_MODIFIER(fallback)
236+
230237
// Modifiers for the 'num_tasks' clause.
231238
OPENMP_NUMTASKS_MODIFIER(strict)
232239

@@ -245,6 +252,7 @@ OPENMP_DOACROSS_MODIFIER(source_omp_cur_iteration)
245252

246253
#undef OPENMP_NUMTASKS_MODIFIER
247254
#undef OPENMP_NUMTHREADS_MODIFIER
255+
#undef OPENMP_DYN_GROUPPRIVATE_MODIFIER
248256
#undef OPENMP_GRAINSIZE_MODIFIER
249257
#undef OPENMP_BIND_KIND
250258
#undef OPENMP_ADJUST_ARGS_KIND

‎clang/include/clang/Basic/OpenMPKinds.h

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -217,6 +217,12 @@ enum OpenMPGrainsizeClauseModifier {
217217
OMPC_GRAINSIZE_unknown
218218
};
219219

220+
enum OpenMPDynGroupprivateClauseModifier {
221+
#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name) OMPC_DYN_GROUPPRIVATE_##Name,
222+
#include "clang/Basic/OpenMPKinds.def"
223+
OMPC_DYN_GROUPPRIVATE_unknown
224+
};
225+
220226
enum OpenMPNumTasksClauseModifier {
221227
#define OPENMP_NUMTASKS_MODIFIER(Name) OMPC_NUMTASKS_##Name,
222228
#include "clang/Basic/OpenMPKinds.def"

‎clang/include/clang/Sema/SemaOpenMP.h

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1385,6 +1385,13 @@ class SemaOpenMP : public SemaBase {
13851385
SourceLocation LParenLoc,
13861386
SourceLocation EndLoc);
13871387

1388+
/// Called on a well-formed 'dyn_groupprivate' clause.
1389+
OMPClause *
1390+
ActOnOpenMPDynGroupprivateClause(OpenMPDynGroupprivateClauseModifier Modifier,
1391+
Expr *Size, SourceLocation StartLoc,
1392+
SourceLocation LParenLoc,
1393+
SourceLocation ModifierLoc, SourceLocation EndLoc);
1394+
13881395
/// Called on well-formed 'doacross' clause.
13891396
OMPClause *
13901397
ActOnOpenMPDoacrossClause(OpenMPDoacrossClauseModifier DepType,

‎clang/lib/AST/OpenMPClause.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -104,6 +104,8 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
104104
return static_cast<const OMPFilterClause *>(C);
105105
case OMPC_ompx_dyn_cgroup_mem:
106106
return static_cast<const OMPXDynCGroupMemClause *>(C);
107+
case OMPC_dyn_groupprivate:
108+
return static_cast<const OMPDynGroupprivateClause *>(C);
107109
case OMPC_default:
108110
case OMPC_proc_bind:
109111
case OMPC_safelen:
@@ -2725,6 +2727,18 @@ void OMPClausePrinter::VisitOMPXDynCGroupMemClause(
27252727
OS << ")";
27262728
}
27272729

2730+
void OMPClausePrinter::VisitOMPDynGroupprivateClause(
2731+
OMPDynGroupprivateClause *Node) {
2732+
OS << "dyn_groupprivate(";
2733+
OpenMPDynGroupprivateClauseModifier Modifier = Node->getModifier();
2734+
if (Modifier != OMPC_DYN_GROUPPRIVATE_unknown) {
2735+
OS << getOpenMPSimpleClauseTypeName(Node->getClauseKind(), Modifier)
2736+
<< ": ";
2737+
}
2738+
Node->getSize()->printPretty(OS, nullptr, Policy, 0);
2739+
OS << ")";
2740+
}
2741+
27282742
void OMPClausePrinter::VisitOMPDoacrossClause(OMPDoacrossClause *Node) {
27292743
OS << "doacross(";
27302744
OpenMPDoacrossClauseModifier DepType = Node->getDependenceType();

‎clang/lib/AST/StmtProfile.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -957,6 +957,12 @@ void OMPClauseProfiler::VisitOMPXDynCGroupMemClause(
957957
if (Expr *Size = C->getSize())
958958
Profiler->VisitStmt(Size);
959959
}
960+
void OMPClauseProfiler::VisitOMPDynGroupprivateClause(
961+
const OMPDynGroupprivateClause *C) {
962+
VistOMPClauseWithPreInit(C);
963+
if (Expr *Size = C->getSize())
964+
Profiler->VisitStmt(Size);
965+
}
960966
void OMPClauseProfiler::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
961967
VisitOMPClauseList(C);
962968
}

‎clang/lib/Basic/OpenMPKinds.cpp

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -171,6 +171,12 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
171171
return OMPC_GRAINSIZE_unknown;
172172
return Type;
173173
}
174+
case OMPC_dyn_groupprivate: {
175+
return llvm::StringSwitch<unsigned>(Str)
176+
#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name) .Case(#Name, OMPC_DYN_GROUPPRIVATE_##Name)
177+
#include "clang/Basic/OpenMPKinds.def"
178+
.Default(OMPC_DYN_GROUPPRIVATE_unknown);
179+
}
174180
case OMPC_num_tasks: {
175181
unsigned Type = llvm::StringSwitch<unsigned>(Str)
176182
#define OPENMP_NUMTASKS_MODIFIER(Name) .Case(#Name, OMPC_NUMTASKS_##Name)
@@ -508,6 +514,16 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
508514
#include "clang/Basic/OpenMPKinds.def"
509515
}
510516
llvm_unreachable("Invalid OpenMP 'grainsize' clause modifier");
517+
case OMPC_dyn_groupprivate:
518+
switch (Type) {
519+
case OMPC_DYN_GROUPPRIVATE_unknown:
520+
return "unknown";
521+
#define OPENMP_DYN_GROUPPRIVATE_MODIFIER(Name) \
522+
case OMPC_DYN_GROUPPRIVATE_##Name: \
523+
return #Name;
524+
#include "clang/Basic/OpenMPKinds.def"
525+
}
526+
llvm_unreachable("Invalid OpenMP 'dyn_groupprivate' clause modifier");
511527
case OMPC_num_tasks:
512528
switch (Type) {
513529
case OMPC_NUMTASKS_unknown:

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 23 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -9489,18 +9489,27 @@ static llvm::Value *emitDeviceID(
94899489
return DeviceID;
94909490
}
94919491

9492-
static llvm::Value *emitDynCGGroupMem(const OMPExecutableDirective &D,
9493-
CodeGenFunction &CGF) {
9494-
llvm::Value *DynCGroupMem = CGF.Builder.getInt32(0);
9495-
9496-
if (auto *DynMemClause = D.getSingleClause<OMPXDynCGroupMemClause>()) {
9497-
CodeGenFunction::RunCleanupsScope DynCGroupMemScope(CGF);
9498-
llvm::Value *DynCGroupMemVal = CGF.EmitScalarExpr(
9499-
DynMemClause->getSize(), /*IgnoreResultAssign=*/true);
9500-
DynCGroupMem = CGF.Builder.CreateIntCast(DynCGroupMemVal, CGF.Int32Ty,
9501-
/*isSigned=*/false);
9502-
}
9503-
return DynCGroupMem;
9492+
static std::pair<llvm::Value *, bool> emitDynCGroupMem(const OMPExecutableDirective &D,
9493+
CodeGenFunction &CGF) {
9494+
llvm::Value *DynGP = CGF.Builder.getInt32(0);
9495+
bool DynGPFallback = false;
9496+
9497+
if (auto *DynGPClause = D.getSingleClause<OMPDynGroupprivateClause>()) {
9498+
CodeGenFunction::RunCleanupsScope DynGPScope(CGF);
9499+
llvm::Value *DynGPVal = CGF.EmitScalarExpr(
9500+
DynGPClause->getSize(), /*IgnoreResultAssign=*/true);
9501+
DynGP = CGF.Builder.CreateIntCast(DynGPVal, CGF.Int32Ty,
9502+
/*isSigned=*/false);
9503+
DynGPFallback = (DynGPClause->getModifier() == OMPC_DYN_GROUPPRIVATE_fallback ||
9504+
OMPC_DYN_GROUPPRIVATE_unknown);
9505+
} else if (auto *OMPXDynCGClause = D.getSingleClause<OMPXDynCGroupMemClause>()) {
9506+
CodeGenFunction::RunCleanupsScope DynCGMemScope(CGF);
9507+
llvm::Value *DynCGMemVal = CGF.EmitScalarExpr(
9508+
OMPXDynCGClause->getSize(), /*IgnoreResultAssign=*/true);
9509+
DynGP = CGF.Builder.CreateIntCast(DynCGMemVal, CGF.Int32Ty,
9510+
/*isSigned=*/false);
9511+
}
9512+
return { DynGP, DynGPFallback };
95049513
}
95059514
static void genMapInfoForCaptures(
95069515
MappableExprsHandler &MEHandler, CodeGenFunction &CGF,
@@ -9710,7 +9719,7 @@ static void emitTargetCallKernelLaunch(
97109719
llvm::Value *RTLoc = OMPRuntime->emitUpdateLocation(CGF, D.getBeginLoc());
97119720
llvm::Value *NumIterations =
97129721
OMPRuntime->emitTargetNumIterationsCall(CGF, D, SizeEmitter);
9713-
llvm::Value *DynCGGroupMem = emitDynCGGroupMem(D, CGF);
9722+
auto [DynCGroupMem, DynCGroupMemFallback] = emitDynCGroupMem(D, CGF);
97149723
llvm::OpenMPIRBuilder::InsertPointTy AllocaIP(
97159724
CGF.AllocaInsertPt->getParent(), CGF.AllocaInsertPt->getIterator());
97169725

@@ -9720,7 +9729,7 @@ static void emitTargetCallKernelLaunch(
97209729

97219730
llvm::OpenMPIRBuilder::TargetKernelArgs Args(
97229731
NumTargetItems, RTArgs, NumIterations, NumTeams, NumThreads,
9723-
DynCGGroupMem, HasNoWait);
9732+
DynCGroupMem, HasNoWait, DynCGroupMemFallback);
97249733

97259734
llvm::OpenMPIRBuilder::InsertPointTy AfterIP =
97269735
cantFail(OMPRuntime->getOMPBuilder().emitKernelLaunch(

‎clang/lib/Parse/ParseOpenMP.cpp

Lines changed: 35 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3039,6 +3039,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
30393039
case OMPC_align:
30403040
case OMPC_message:
30413041
case OMPC_ompx_dyn_cgroup_mem:
3042+
case OMPC_dyn_groupprivate:
30423043
// OpenMP [2.5, Restrictions]
30433044
// At most one num_threads clause can appear on the directive.
30443045
// OpenMP [2.8.1, simd construct, Restrictions]
@@ -3077,7 +3078,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
30773078
PP.LookAhead(/*N=*/0).isNot(tok::l_paren))
30783079
Clause = ParseOpenMPClause(CKind, WrongDirective);
30793080
else if (CKind == OMPC_grainsize || CKind == OMPC_num_tasks ||
3080-
CKind == OMPC_num_threads)
3081+
CKind == OMPC_num_threads || CKind == OMPC_dyn_groupprivate)
30813082
Clause = ParseOpenMPSingleExprWithArgClause(DKind, CKind, WrongDirective);
30823083
else
30833084
Clause = ParseOpenMPSingleExprClause(CKind, WrongDirective);
@@ -3835,6 +3836,37 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
38353836
Arg.push_back(OMPC_GRAINSIZE_unknown);
38363837
KLoc.emplace_back();
38373838
}
3839+
} else if (Kind == OMPC_dyn_groupprivate) {
3840+
// Parse optional <dyn_groupprivate modifier> ':'
3841+
OpenMPDynGroupprivateClauseModifier Modifier =
3842+
static_cast<OpenMPDynGroupprivateClauseModifier>(getOpenMPSimpleClauseType(
3843+
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
3844+
getLangOpts()));
3845+
if (getLangOpts().OpenMP >= 51) {
3846+
if (NextToken().is(tok::colon)) {
3847+
Arg.push_back(Modifier);
3848+
KLoc.push_back(Tok.getLocation());
3849+
// Parse modifier
3850+
ConsumeAnyToken();
3851+
// Parse ':'
3852+
ConsumeAnyToken();
3853+
} else {
3854+
if (Modifier == OMPC_DYN_GROUPPRIVATE_strict) {
3855+
Diag(Tok, diag::err_modifier_expected_colon) << "strict";
3856+
// Parse modifier
3857+
ConsumeAnyToken();
3858+
} else if (Modifier == OMPC_DYN_GROUPPRIVATE_fallback) {
3859+
Diag(Tok, diag::err_modifier_expected_colon) << "fallback";
3860+
// Parse modifier
3861+
ConsumeAnyToken();
3862+
}
3863+
Arg.push_back(OMPC_DYN_GROUPPRIVATE_unknown);
3864+
KLoc.emplace_back();
3865+
}
3866+
} else {
3867+
Arg.push_back(OMPC_DYN_GROUPPRIVATE_unknown);
3868+
KLoc.emplace_back();
3869+
}
38383870
} else if (Kind == OMPC_num_tasks) {
38393871
// Parse optional <num_tasks modifier> ':'
38403872
OpenMPNumTasksClauseModifier Modifier =
@@ -3913,7 +3945,8 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
39133945
(Kind == OMPC_dist_schedule && DelimLoc.isValid()) ||
39143946
Kind == OMPC_if || Kind == OMPC_device ||
39153947
Kind == OMPC_grainsize || Kind == OMPC_num_tasks ||
3916-
Kind == OMPC_num_threads;
3948+
Kind == OMPC_num_threads ||
3949+
Kind == OMPC_dyn_groupprivate;
39173950
if (NeedAnExpression) {
39183951
SourceLocation ELoc = Tok.getLocation();
39193952
ExprResult LHS(

0 commit comments

Comments
 (0)