Skip to content

Commit b146a1a

Browse files
author
Chandra Ghale
committed
Few more fixes with ref from spec
1 parent 4c36ba7 commit b146a1a

7 files changed

+136
-160
lines changed

clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 47 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -4915,12 +4915,12 @@ void CGOpenMPRuntime::emitPrivateReduction(
49154915
// - Thread enters critical section.
49164916
// - Reads its private value from LHSExprs[i].
49174917
// - Updates __shared_reduction_var[i] = RedOp_i(__shared_reduction_var[i],
4918-
// LHSExprs[i]).
4918+
// Privates[i]).
49194919
// - Exits critical section.
49204920
//
49214921
// Call __kmpc_barrier after combining.
49224922
//
4923-
// Each thread copies __shared_reduction_var[i] back to LHSExprs[i].
4923+
// Each thread copies __shared_reduction_var[i] back to RHSExprs[i].
49244924
//
49254925
// Final __kmpc_barrier to synchronize after broadcasting
49264926
QualType PrivateType = Privates->getType();
@@ -5025,7 +5025,7 @@ void CGOpenMPRuntime::emitPrivateReduction(
50255025
const Expr *ReductionOp = ReductionOps;
50265026
const OMPDeclareReductionDecl *CurrentUDR = getReductionInit(ReductionOp);
50275027
LValue SharedLV = CGF.MakeAddrLValue(SharedResult, PrivateType);
5028-
LValue LHSLV = CGF.EmitLValue(LHSExprs);
5028+
LValue LHSLV = CGF.EmitLValue(Privates);
50295029

50305030
auto EmitCriticalReduction = [&](auto ReductionGen) {
50315031
std::string CriticalName = getName({"reduction_critical"});
@@ -5114,7 +5114,7 @@ void CGOpenMPRuntime::emitPrivateReduction(
51145114
else
51155115
FinalResultVal = CGF.EmitLoadOfScalar(SharedLV1, Loc);
51165116

5117-
LValue TargetLHSLV = CGF.EmitLValue(LHSExprs);
5117+
LValue TargetLHSLV = CGF.EmitLValue(RHSExprs);
51185118
if (IsAggregate) {
51195119
CGF.EmitAggregateCopy(TargetLHSLV,
51205120
CGF.MakeAddrLValue(FinalResultAddr, PrivateType),
@@ -5126,13 +5126,23 @@ void CGOpenMPRuntime::emitPrivateReduction(
51265126
CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
51275127
CGM.getModule(), OMPRTL___kmpc_barrier),
51285128
BarrierArgs);
5129+
5130+
// Combiner with original list item
5131+
auto OriginalListCombiner = [&](CodeGenFunction &CGF,
5132+
PrePostActionTy &Action) {
5133+
Action.Enter(CGF);
5134+
emitSingleReductionCombiner(CGF, ReductionOps, Privates,
5135+
cast<DeclRefExpr>(LHSExprs),
5136+
cast<DeclRefExpr>(RHSExprs));
5137+
};
5138+
EmitCriticalReduction(OriginalListCombiner);
51295139
}
51305140

51315141
void CGOpenMPRuntime::emitReduction(CodeGenFunction &CGF, SourceLocation Loc,
5132-
ArrayRef<const Expr *> Privates,
5133-
ArrayRef<const Expr *> LHSExprs,
5134-
ArrayRef<const Expr *> RHSExprs,
5135-
ArrayRef<const Expr *> ReductionOps,
5142+
ArrayRef<const Expr *> OrgPrivates,
5143+
ArrayRef<const Expr *> OrgLHSExprs,
5144+
ArrayRef<const Expr *> OrgRHSExprs,
5145+
ArrayRef<const Expr *> OrgReductionOps,
51365146
ReductionOptionsTy Options) {
51375147
if (!CGF.HaveInsertPoint())
51385148
return;
@@ -5179,10 +5189,10 @@ void CGOpenMPRuntime::emitReduction(CodeGenFunction &CGF, SourceLocation Loc,
51795189

51805190
if (SimpleReduction) {
51815191
CodeGenFunction::RunCleanupsScope Scope(CGF);
5182-
const auto *IPriv = Privates.begin();
5183-
const auto *ILHS = LHSExprs.begin();
5184-
const auto *IRHS = RHSExprs.begin();
5185-
for (const Expr *E : ReductionOps) {
5192+
const auto *IPriv = OrgPrivates.begin();
5193+
const auto *ILHS = OrgLHSExprs.begin();
5194+
const auto *IRHS = OrgRHSExprs.begin();
5195+
for (const Expr *E : OrgReductionOps) {
51865196
emitSingleReductionCombiner(CGF, E, *IPriv, cast<DeclRefExpr>(*ILHS),
51875197
cast<DeclRefExpr>(*IRHS));
51885198
++IPriv;
@@ -5192,6 +5202,26 @@ void CGOpenMPRuntime::emitReduction(CodeGenFunction &CGF, SourceLocation Loc,
51925202
return;
51935203
}
51945204

5205+
// Filter out shared reduction variables based on IsPrivateVarReduction flag.
5206+
// Only keep entries where the corresponding variable is not private.
5207+
SmallVector<const Expr *> FilteredPrivates, FilteredLHSExprs,
5208+
FilteredRHSExprs, FilteredReductionOps;
5209+
for (unsigned I : llvm::seq<unsigned>(
5210+
std::min(OrgReductionOps.size(), OrgLHSExprs.size()))) {
5211+
if (!Options.IsPrivateVarReduction[I]) {
5212+
FilteredPrivates.emplace_back(OrgPrivates[I]);
5213+
FilteredLHSExprs.emplace_back(OrgLHSExprs[I]);
5214+
FilteredRHSExprs.emplace_back(OrgRHSExprs[I]);
5215+
FilteredReductionOps.emplace_back(OrgReductionOps[I]);
5216+
}
5217+
}
5218+
// Wrap filtered vectors in ArrayRef for downstream shared reduction
5219+
// processing.
5220+
ArrayRef<const Expr *> Privates = FilteredPrivates;
5221+
ArrayRef<const Expr *> LHSExprs = FilteredLHSExprs;
5222+
ArrayRef<const Expr *> RHSExprs = FilteredRHSExprs;
5223+
ArrayRef<const Expr *> ReductionOps = FilteredReductionOps;
5224+
51955225
// 1. Build a list of reduction variables.
51965226
// void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
51975227
auto Size = RHSExprs.size();
@@ -5439,11 +5469,12 @@ void CGOpenMPRuntime::emitReduction(CodeGenFunction &CGF, SourceLocation Loc,
54395469
"PrivateVarReduction: ReductionOps size mismatch");
54405470
assert(LHSExprs.size() == Options.IsPrivateVarReduction.size() &&
54415471
"PrivateVarReduction: IsPrivateVarReduction size mismatch");
5442-
for (unsigned I :
5443-
llvm::seq<unsigned>(std::min(ReductionOps.size(), LHSExprs.size()))) {
5472+
5473+
for (unsigned I : llvm::seq<unsigned>(
5474+
std::min(OrgReductionOps.size(), OrgLHSExprs.size()))) {
54445475
if (Options.IsPrivateVarReduction[I])
5445-
emitPrivateReduction(CGF, Loc, Privates[I], LHSExprs[I], RHSExprs[I],
5446-
ReductionOps[I]);
5476+
emitPrivateReduction(CGF, Loc, OrgPrivates[I], OrgLHSExprs[I],
5477+
OrgRHSExprs[I], OrgReductionOps[I]);
54475478
}
54485479
}
54495480

clang/lib/Sema/SemaOpenMP.cpp

Lines changed: 19 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -19015,34 +19015,14 @@ static bool actOnOMPReductionKindClause(
1901519015
reportOriginalDsa(S, Stack, D, DVar);
1901619016
continue;
1901719017
}
19018-
// OpenMP 6.0 [ 7.6.10 ]
19019-
// Support Reduction over private variables with reduction clause.
19020-
// A list item in a reduction clause can now be private in the enclosing
19021-
// context. For orphaned constructs it is assumed to be shared unless the
19022-
// original(private) modifier appears in the clause.
19023-
DVar = Stack->getImplicitDSA(D, true);
19024-
bool IsOrphaned = false;
19025-
OpenMPDirectiveKind CurrDir = Stack->getCurrentDirective();
19026-
OpenMPDirectiveKind ParentDir = Stack->getParentDirective();
19027-
// Check if the construct is orphaned (has no enclosing OpenMP context)
19028-
IsOrphaned = ParentDir == OMPD_unknown;
19029-
// OpenMP 6.0: Private DSA check
19030-
IsPrivate =
19031-
(S.getLangOpts().OpenMP > 52) &&
19032-
((isOpenMPPrivate(DVar.CKind) && DVar.CKind != OMPC_reduction &&
19033-
isOpenMPWorksharingDirective(CurrDir) &&
19034-
!isOpenMPParallelDirective(CurrDir) &&
19035-
!isOpenMPTeamsDirective(CurrDir) &&
19036-
!isOpenMPSimdDirective(ParentDir)) ||
19037-
(IsOrphaned && DVar.CKind == OMPC_unknown) ||
19038-
RD.OrigSharingModifier != OMPC_ORIGINAL_SHARING_shared);
1903919018

1904019019
// OpenMP [2.14.3.6, Restrictions, p.1]
1904119020
// A list item that appears in a reduction clause of a worksharing
1904219021
// construct must be shared in the parallel regions to which any of the
1904319022
// worksharing regions arising from the worksharing construct bind.
1904419023

19045-
if (!IsPrivate && isOpenMPWorksharingDirective(CurrDir) &&
19024+
if (S.getLangOpts().OpenMP <= 52 &&
19025+
isOpenMPWorksharingDirective(CurrDir) &&
1904619026
!isOpenMPParallelDirective(CurrDir) &&
1904719027
!isOpenMPTeamsDirective(CurrDir)) {
1904819028
DVar = Stack->getImplicitDSA(D, true);
@@ -19053,6 +19033,23 @@ static bool actOnOMPReductionKindClause(
1905319033
reportOriginalDsa(S, Stack, D, DVar);
1905419034
continue;
1905519035
}
19036+
} else if (isOpenMPWorksharingDirective(CurrDir) &&
19037+
!isOpenMPParallelDirective(CurrDir) &&
19038+
!isOpenMPTeamsDirective(CurrDir)) {
19039+
// OpenMP 6.0 [ 7.6.10 ]
19040+
// Support Reduction over private variables with reduction clause.
19041+
// A list item in a reduction clause can now be private in the enclosing
19042+
// context. For orphaned constructs it is assumed to be shared unless
19043+
// the original(private) modifier appears in the clause.
19044+
DVar = Stack->getImplicitDSA(D, true);
19045+
// Determine if the variable should be considered private
19046+
IsPrivate = DVar.CKind != OMPC_shared;
19047+
bool IsOrphaned = false;
19048+
OpenMPDirectiveKind ParentDir = Stack->getParentDirective();
19049+
IsOrphaned = ParentDir == OMPD_unknown;
19050+
if ((IsOrphaned &&
19051+
RD.OrigSharingModifier == OMPC_ORIGINAL_SHARING_private))
19052+
IsPrivate = true;
1905619053
}
1905719054
} else {
1905819055
// Threadprivates cannot be shared between threads, so dignose if the base

clang/test/OpenMP/distribute_simd_misc_messages.c

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -508,6 +508,7 @@ void test_collapse(void) {
508508
#pragma omp distribute simd collapse(5 - 5)
509509
for (i = 0; i < 16; ++i)
510510
;
511+
#if defined(_OPENMP) && (_OPENMP <= 202111)
511512
// expected-note@+3 2 {{defined as reduction}}
512513
#pragma omp target
513514
#pragma omp teams
@@ -520,7 +521,7 @@ void test_collapse(void) {
520521
#pragma omp for reduction(+ : i, j)
521522
for (int k = 0; k < 16; ++k)
522523
i += j;
523-
524+
#endif
524525
#pragma omp target
525526
#pragma omp teams
526527
for (i = 0; i < 16; ++i)

0 commit comments

Comments
 (0)