r327985 - [OPENMP, NVPTX] Globalization of the private redeclarations.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Tue Mar 20 07:45:59 PDT 2018
Author: abataev
Date: Tue Mar 20 07:45:59 2018
New Revision: 327985
URL: http://llvm.org/viewvc/llvm-project?rev=327985&view=rev
Log:
[OPENMP, NVPTX] Globalization of the private redeclarations.
If the generic codegen is enabled and private copy of the original
variable escapes the declaration context, this private copy should be
globalized just like it was the original variable.
Modified:
cfe/trunk/include/clang/Basic/Attr.td
cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/test/OpenMP/nvptx_data_sharing.cpp
Modified: cfe/trunk/include/clang/Basic/Attr.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=327985&r1=327984&r2=327985&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Tue Mar 20 07:45:59 2018
@@ -2779,6 +2779,14 @@ def OMPCaptureKind : Attr {
let Documentation = [Undocumented];
}
+def OMPReferencedVar : Attr {
+ // This attribute has no spellings as it is only ever created implicitly.
+ let Spellings = [];
+ let SemaHandler = 0;
+ let Args = [ExprArgument<"Ref">];
+ let Documentation = [Undocumented];
+}
+
def OMPDeclareSimdDecl : Attr {
let Spellings = [Pragma<"omp", "declare simd">];
let Subjects = SubjectList<[Function]>;
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp?rev=327985&r1=327984&r2=327985&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp Tue Mar 20 07:45:59 2018
@@ -171,36 +171,48 @@ class CheckVarsEscapingDeclContext final
: public ConstStmtVisitor<CheckVarsEscapingDeclContext> {
CodeGenFunction &CGF;
llvm::SetVector<const ValueDecl *> EscapedDecls;
+ llvm::SetVector<const ValueDecl *> EscapedVariableLengthDecls;
llvm::SmallPtrSet<const Decl *, 4> EscapedParameters;
- llvm::SmallPtrSet<const ValueDecl *, 4> IgnoredDecls;
bool AllEscaped = false;
RecordDecl *GlobalizedRD = nullptr;
llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
void markAsEscaped(const ValueDecl *VD) {
- if (IgnoredDecls.count(VD))
- return;
+ VD = cast<ValueDecl>(VD->getCanonicalDecl());
// Variables captured by value must be globalized.
if (auto *CSI = CGF.CapturedStmtInfo) {
if (const FieldDecl *FD = CSI->lookup(cast<VarDecl>(VD))) {
+ if (!FD->hasAttrs())
+ return;
+ const auto *Attr = FD->getAttr<OMPCaptureKindAttr>();
+ if (!Attr)
+ return;
+ if (!isOpenMPPrivate(
+ static_cast<OpenMPClauseKind>(Attr->getCaptureKind())) ||
+ Attr->getCaptureKind() == OMPC_map)
+ return;
if (FD->getType()->isReferenceType())
return;
+ assert(!VD->getType()->isVariablyModifiedType() &&
+ "Parameter captured by value with variably modified type");
EscapedParameters.insert(VD);
}
}
- EscapedDecls.insert(VD);
+ if (VD->getType()->isVariablyModifiedType())
+ EscapedVariableLengthDecls.insert(VD);
+ else
+ EscapedDecls.insert(VD);
}
void VisitValueDecl(const ValueDecl *VD) {
- if (VD->getType()->isLValueReferenceType()) {
+ if (VD->getType()->isLValueReferenceType())
markAsEscaped(VD);
- if (const auto *VarD = dyn_cast<VarDecl>(VD)) {
- if (!isa<ParmVarDecl>(VarD) && VarD->hasInit()) {
- const bool SavedAllEscaped = AllEscaped;
- AllEscaped = true;
- Visit(VarD->getInit());
- AllEscaped = SavedAllEscaped;
- }
+ if (const auto *VarD = dyn_cast<VarDecl>(VD)) {
+ if (!isa<ParmVarDecl>(VarD) && VarD->hasInit()) {
+ const bool SavedAllEscaped = AllEscaped;
+ AllEscaped = VD->getType()->isLValueReferenceType();
+ Visit(VarD->getInit());
+ AllEscaped = SavedAllEscaped;
}
}
}
@@ -265,9 +277,7 @@ class CheckVarsEscapingDeclContext final
}
public:
- CheckVarsEscapingDeclContext(CodeGenFunction &CGF,
- ArrayRef<const ValueDecl *> IgnoredDecls)
- : CGF(CGF), IgnoredDecls(IgnoredDecls.begin(), IgnoredDecls.end()) {}
+ CheckVarsEscapingDeclContext(CodeGenFunction &CGF) : CGF(CGF) {}
virtual ~CheckVarsEscapingDeclContext() = default;
void VisitDeclStmt(const DeclStmt *S) {
if (!S)
@@ -420,6 +430,12 @@ public:
const llvm::SmallPtrSetImpl<const Decl *> &getEscapedParameters() const {
return EscapedParameters;
}
+
+ /// Returns the list of the escaped variables with the variably modified
+ /// types.
+ ArrayRef<const ValueDecl *> getEscapedVariableLengthDecls() const {
+ return EscapedVariableLengthDecls.getArrayRef();
+ }
};
} // anonymous namespace
@@ -1247,63 +1263,103 @@ void CGOpenMPRuntimeNVPTX::emitGenericVa
const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
if (I == FunctionGlobalizedDecls.end())
return;
- const RecordDecl *GlobalizedVarsRecord = I->getSecond().GlobalRecord;
- QualType RecTy = CGM.getContext().getRecordType(GlobalizedVarsRecord);
+ if (const RecordDecl *GlobalizedVarsRecord = I->getSecond().GlobalRecord) {
+ QualType RecTy = CGM.getContext().getRecordType(GlobalizedVarsRecord);
- // Recover pointer to this function's global record. The runtime will
- // handle the specifics of the allocation of the memory.
- // Use actual memory size of the record including the padding
- // for alignment purposes.
- unsigned Alignment =
- CGM.getContext().getTypeAlignInChars(RecTy).getQuantity();
- unsigned GlobalRecordSize =
- CGM.getContext().getTypeSizeInChars(RecTy).getQuantity();
- GlobalRecordSize = llvm::alignTo(GlobalRecordSize, Alignment);
- // TODO: allow the usage of shared memory to be controlled by
- // the user, for now, default to global.
- llvm::Value *GlobalRecordSizeArg[] = {
- llvm::ConstantInt::get(CGM.SizeTy, GlobalRecordSize),
- CGF.Builder.getInt16(/*UseSharedMemory=*/0)};
- llvm::Value *GlobalRecValue = CGF.EmitRuntimeCall(
- createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_data_sharing_push_stack),
- GlobalRecordSizeArg);
- llvm::Value *GlobalRecCastAddr = Bld.CreatePointerBitCastOrAddrSpaceCast(
- GlobalRecValue, CGF.ConvertTypeForMem(RecTy)->getPointerTo());
- LValue Base = CGF.MakeNaturalAlignPointeeAddrLValue(GlobalRecCastAddr, RecTy);
- I->getSecond().GlobalRecordAddr = GlobalRecValue;
-
- // Emit the "global alloca" which is a GEP from the global declaration record
- // using the pointer returned by the runtime.
- for (auto &Rec : I->getSecond().LocalVarData) {
- bool EscapedParam = I->getSecond().EscapedParameters.count(Rec.first);
- llvm::Value *ParValue;
- if (EscapedParam) {
- const auto *VD = cast<VarDecl>(Rec.first);
- LValue ParLVal =
- CGF.MakeAddrLValue(CGF.GetAddrOfLocalVar(VD), VD->getType());
- ParValue = CGF.EmitLoadOfScalar(ParLVal, Loc);
- }
- const FieldDecl *FD = Rec.second.first;
- LValue VarAddr = CGF.EmitLValueForField(Base, FD);
- Rec.second.second = VarAddr.getAddress();
- if (EscapedParam) {
- const auto *VD = cast<VarDecl>(Rec.first);
- CGF.EmitStoreOfScalar(ParValue, VarAddr);
- I->getSecond().MappedParams->setVarAddr(CGF, VD, VarAddr.getAddress());
+ // Recover pointer to this function's global record. The runtime will
+ // handle the specifics of the allocation of the memory.
+ // Use actual memory size of the record including the padding
+ // for alignment purposes.
+ unsigned Alignment =
+ CGM.getContext().getTypeAlignInChars(RecTy).getQuantity();
+ unsigned GlobalRecordSize =
+ CGM.getContext().getTypeSizeInChars(RecTy).getQuantity();
+ GlobalRecordSize = llvm::alignTo(GlobalRecordSize, Alignment);
+ // TODO: allow the usage of shared memory to be controlled by
+ // the user, for now, default to global.
+ llvm::Value *GlobalRecordSizeArg[] = {
+ llvm::ConstantInt::get(CGM.SizeTy, GlobalRecordSize),
+ CGF.Builder.getInt16(/*UseSharedMemory=*/0)};
+ llvm::Value *GlobalRecValue = CGF.EmitRuntimeCall(
+ createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_data_sharing_push_stack),
+ GlobalRecordSizeArg);
+ llvm::Value *GlobalRecCastAddr = Bld.CreatePointerBitCastOrAddrSpaceCast(
+ GlobalRecValue, CGF.ConvertTypeForMem(RecTy)->getPointerTo());
+ LValue Base =
+ CGF.MakeNaturalAlignPointeeAddrLValue(GlobalRecCastAddr, RecTy);
+ I->getSecond().GlobalRecordAddr = GlobalRecValue;
+
+ // Emit the "global alloca" which is a GEP from the global declaration
+ // record using the pointer returned by the runtime.
+ for (auto &Rec : I->getSecond().LocalVarData) {
+ bool EscapedParam = I->getSecond().EscapedParameters.count(Rec.first);
+ llvm::Value *ParValue;
+ if (EscapedParam) {
+ const auto *VD = cast<VarDecl>(Rec.first);
+ LValue ParLVal =
+ CGF.MakeAddrLValue(CGF.GetAddrOfLocalVar(VD), VD->getType());
+ ParValue = CGF.EmitLoadOfScalar(ParLVal, Loc);
+ }
+ const FieldDecl *FD = Rec.second.first;
+ LValue VarAddr = CGF.EmitLValueForField(Base, FD);
+ Rec.second.second = VarAddr.getAddress();
+ if (EscapedParam) {
+ const auto *VD = cast<VarDecl>(Rec.first);
+ CGF.EmitStoreOfScalar(ParValue, VarAddr);
+ I->getSecond().MappedParams->setVarAddr(CGF, VD, VarAddr.getAddress());
+ }
}
}
+ for (const ValueDecl *VD : I->getSecond().EscapedVariableLengthDecls) {
+ // Recover pointer to this function's global record. The runtime will
+ // handle the specifics of the allocation of the memory.
+ // Use actual memory size of the record including the padding
+ // for alignment purposes.
+ auto &Bld = CGF.Builder;
+ llvm::Value *Size = CGF.getTypeSize(VD->getType());
+ CharUnits Align = CGM.getContext().getDeclAlign(VD);
+ Size = Bld.CreateNUWAdd(
+ Size, llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity() - 1));
+ llvm::Value *AlignVal =
+ llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity());
+ Size = Bld.CreateUDiv(Size, AlignVal);
+ Size = Bld.CreateNUWMul(Size, AlignVal);
+ // TODO: allow the usage of shared memory to be controlled by
+ // the user, for now, default to global.
+ llvm::Value *GlobalRecordSizeArg[] = {
+ Size, CGF.Builder.getInt16(/*UseSharedMemory=*/0)};
+ llvm::Value *GlobalRecValue = CGF.EmitRuntimeCall(
+ createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_data_sharing_push_stack),
+ GlobalRecordSizeArg);
+ llvm::Value *GlobalRecCastAddr = Bld.CreatePointerBitCastOrAddrSpaceCast(
+ GlobalRecValue, CGF.ConvertTypeForMem(VD->getType())->getPointerTo());
+ LValue Base = CGF.MakeAddrLValue(GlobalRecCastAddr, VD->getType(),
+ CGM.getContext().getDeclAlign(VD),
+ AlignmentSource::Decl);
+ I->getSecond().MappedParams->setVarAddr(CGF, cast<VarDecl>(VD),
+ Base.getAddress());
+ I->getSecond().EscapedVariableLengthDeclsAddrs.emplace_back(GlobalRecValue);
+ }
I->getSecond().MappedParams->apply(CGF);
}
void CGOpenMPRuntimeNVPTX::emitGenericVarsEpilog(CodeGenFunction &CGF) {
const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
- if (I != FunctionGlobalizedDecls.end() && I->getSecond().GlobalRecordAddr) {
+ if (I != FunctionGlobalizedDecls.end()) {
I->getSecond().MappedParams->restore(CGF);
if (!CGF.HaveInsertPoint())
return;
- CGF.EmitRuntimeCall(
- createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_data_sharing_pop_stack),
- I->getSecond().GlobalRecordAddr);
+ for (llvm::Value *Addr :
+ llvm::reverse(I->getSecond().EscapedVariableLengthDeclsAddrs)) {
+ CGF.EmitRuntimeCall(
+ createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_data_sharing_pop_stack),
+ Addr);
+ }
+ if (I->getSecond().GlobalRecordAddr) {
+ CGF.EmitRuntimeCall(
+ createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_data_sharing_pop_stack),
+ I->getSecond().GlobalRecordAddr);
+ }
}
}
@@ -2937,7 +2993,6 @@ void CGOpenMPRuntimeNVPTX::emitFunctionP
assert(D && "Expected function or captured|block decl.");
assert(FunctionGlobalizedDecls.count(CGF.CurFn) == 0 &&
"Function is registered already.");
- SmallVector<const ValueDecl *, 4> IgnoredDecls;
const Stmt *Body = nullptr;
bool NeedToDelayGlobalization = false;
if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
@@ -2946,22 +3001,16 @@ void CGOpenMPRuntimeNVPTX::emitFunctionP
Body = BD->getBody();
} else if (const auto *CD = dyn_cast<CapturedDecl>(D)) {
Body = CD->getBody();
- if (CGF.CapturedStmtInfo->getKind() == CR_OpenMP) {
- NeedToDelayGlobalization = true;
- if (const auto *CS = dyn_cast<CapturedStmt>(Body)) {
- IgnoredDecls.reserve(CS->capture_size());
- for (const auto &Capture : CS->captures())
- if (Capture.capturesVariable())
- IgnoredDecls.emplace_back(Capture.getCapturedVar());
- }
- }
+ NeedToDelayGlobalization = CGF.CapturedStmtInfo->getKind() == CR_OpenMP;
}
if (!Body)
return;
- CheckVarsEscapingDeclContext VarChecker(CGF, IgnoredDecls);
+ CheckVarsEscapingDeclContext VarChecker(CGF);
VarChecker.Visit(Body);
const RecordDecl *GlobalizedVarsRecord = VarChecker.getGlobalizedRecord();
- if (!GlobalizedVarsRecord)
+ ArrayRef<const ValueDecl *> EscapedVariableLengthDecls =
+ VarChecker.getEscapedVariableLengthDecls();
+ if (!GlobalizedVarsRecord && EscapedVariableLengthDecls.empty())
return;
auto I = FunctionGlobalizedDecls.try_emplace(CGF.CurFn).first;
I->getSecond().MappedParams =
@@ -2970,8 +3019,11 @@ void CGOpenMPRuntimeNVPTX::emitFunctionP
I->getSecond().EscapedParameters.insert(
VarChecker.getEscapedParameters().begin(),
VarChecker.getEscapedParameters().end());
+ I->getSecond().EscapedVariableLengthDecls.append(
+ EscapedVariableLengthDecls.begin(), EscapedVariableLengthDecls.end());
DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {
+ assert(VD->isCanonicalDecl() && "Expected canonical declaration");
const FieldDecl *FD = VarChecker.getFieldForGlobalizedVar(VD);
Data.insert(std::make_pair(VD, std::make_pair(FD, Address::invalid())));
}
@@ -2991,13 +3043,25 @@ void CGOpenMPRuntimeNVPTX::emitFunctionP
Address CGOpenMPRuntimeNVPTX::getAddressOfLocalVariable(CodeGenFunction &CGF,
const VarDecl *VD) {
+ VD = VD->getCanonicalDecl();
auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
if (I == FunctionGlobalizedDecls.end())
return Address::invalid();
auto VDI = I->getSecond().LocalVarData.find(VD);
- if (VDI == I->getSecond().LocalVarData.end())
- return Address::invalid();
- return VDI->second.second;
+ if (VDI != I->getSecond().LocalVarData.end())
+ return VDI->second.second;
+ if (VD->hasAttrs()) {
+ for (specific_attr_iterator<OMPReferencedVarAttr> IT(VD->attr_begin()),
+ E(VD->attr_end());
+ IT != E; ++IT) {
+ auto VDI = I->getSecond().LocalVarData.find(
+ cast<VarDecl>(cast<DeclRefExpr>(IT->getRef())->getDecl())
+ ->getCanonicalDecl());
+ if (VDI != I->getSecond().LocalVarData.end())
+ return VDI->second.second;
+ }
+ }
+ return Address::invalid();
}
void CGOpenMPRuntimeNVPTX::functionFinished(CodeGenFunction &CGF) {
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h?rev=327985&r1=327984&r2=327985&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h Tue Mar 20 07:45:59 2018
@@ -342,9 +342,11 @@ private:
using EscapedParamsTy = llvm::SmallPtrSet<const Decl *, 4>;
struct FunctionData {
DeclToAddrMapTy LocalVarData;
+ EscapedParamsTy EscapedParameters;
+ llvm::SmallVector<const ValueDecl*, 4> EscapedVariableLengthDecls;
+ llvm::SmallVector<llvm::Value *, 4> EscapedVariableLengthDeclsAddrs;
const RecordDecl *GlobalRecord = nullptr;
llvm::Value *GlobalRecordAddr = nullptr;
- EscapedParamsTy EscapedParameters;
std::unique_ptr<CodeGenFunction::OMPMapVars> MappedParams;
};
/// Maps the function to the list of the globalized variables with their
Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=327985&r1=327984&r2=327985&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Tue Mar 20 07:45:59 2018
@@ -1263,6 +1263,7 @@ static void emitEmptyBoundParameters(Cod
void CodeGenFunction::EmitOMPParallelDirective(const OMPParallelDirective &S) {
// Emit parallel region as a standalone region.
auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
+ Action.Enter(CGF);
OMPPrivateScope PrivateScope(CGF);
bool Copyins = CGF.EmitOMPCopyinClause(S);
(void)CGF.EmitOMPFirstprivateClause(S, PrivateScope);
@@ -1277,7 +1278,6 @@ void CodeGenFunction::EmitOMPParallelDir
CGF.EmitOMPPrivateClause(S, PrivateScope);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.EmitStmt(S.getCapturedStmt(OMPD_parallel)->getCapturedStmt());
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_parallel);
};
@@ -3981,12 +3981,12 @@ static void emitCommonOMPTargetDirective
static void emitTargetRegion(CodeGenFunction &CGF, const OMPTargetDirective &S,
PrePostActionTy &Action) {
+ Action.Enter(CGF);
CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
(void)CGF.EmitOMPFirstprivateClause(S, PrivateScope);
CGF.EmitOMPPrivateClause(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.EmitStmt(S.getCapturedStmt(OMPD_target)->getCapturedStmt());
}
@@ -4039,12 +4039,12 @@ static void emitCommonOMPTeamsDirective(
void CodeGenFunction::EmitOMPTeamsDirective(const OMPTeamsDirective &S) {
// Emit teams region as a standalone region.
auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
+ Action.Enter(CGF);
OMPPrivateScope PrivateScope(CGF);
(void)CGF.EmitOMPFirstprivateClause(S, PrivateScope);
CGF.EmitOMPPrivateClause(S, PrivateScope);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.EmitStmt(S.getCapturedStmt(OMPD_teams)->getCapturedStmt());
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
};
@@ -4059,12 +4059,12 @@ static void emitTargetTeamsRegion(CodeGe
Action.Enter(CGF);
// Emit teams region as a standalone region.
auto &&CodeGen = [&S, CS](CodeGenFunction &CGF, PrePostActionTy &Action) {
+ Action.Enter(CGF);
CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
(void)CGF.EmitOMPFirstprivateClause(S, PrivateScope);
CGF.EmitOMPPrivateClause(S, PrivateScope);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.EmitStmt(CS->getCapturedStmt());
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
};
@@ -4106,10 +4106,10 @@ emitTargetTeamsDistributeRegion(CodeGenF
// Emit teams region as a standalone region.
auto &&CodeGen = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
PrePostActionTy &Action) {
+ Action.Enter(CGF);
CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_distribute,
CodeGenDistribute);
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
@@ -4152,10 +4152,10 @@ static void emitTargetTeamsDistributeSim
// Emit teams region as a standalone region.
auto &&CodeGen = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
PrePostActionTy &Action) {
+ Action.Enter(CGF);
CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_distribute,
CodeGenDistribute);
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
@@ -4197,10 +4197,10 @@ void CodeGenFunction::EmitOMPTeamsDistri
// Emit teams region as a standalone region.
auto &&CodeGen = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
PrePostActionTy &Action) {
+ Action.Enter(CGF);
OMPPrivateScope PrivateScope(CGF);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_distribute,
CodeGenDistribute);
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
@@ -4219,10 +4219,10 @@ void CodeGenFunction::EmitOMPTeamsDistri
// Emit teams region as a standalone region.
auto &&CodeGen = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
PrePostActionTy &Action) {
+ Action.Enter(CGF);
OMPPrivateScope PrivateScope(CGF);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_simd,
CodeGenDistribute);
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
@@ -4242,10 +4242,10 @@ void CodeGenFunction::EmitOMPTeamsDistri
// Emit teams region as a standalone region.
auto &&CodeGen = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
PrePostActionTy &Action) {
+ Action.Enter(CGF);
OMPPrivateScope PrivateScope(CGF);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_distribute,
CodeGenDistribute);
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
@@ -4265,10 +4265,10 @@ void CodeGenFunction::EmitOMPTeamsDistri
// Emit teams region as a standalone region.
auto &&CodeGen = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
PrePostActionTy &Action) {
+ Action.Enter(CGF);
OMPPrivateScope PrivateScope(CGF);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.CGM.getOpenMPRuntime().emitInlinedDirective(
CGF, OMPD_distribute, CodeGenDistribute, /*HasCancel=*/false);
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
@@ -4290,10 +4290,10 @@ static void emitTargetTeamsDistributePar
// Emit teams region as a standalone region.
auto &&CodeGenTeams = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
PrePostActionTy &Action) {
+ Action.Enter(CGF);
CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.CGM.getOpenMPRuntime().emitInlinedDirective(
CGF, OMPD_distribute, CodeGenDistribute, /*HasCancel=*/false);
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
@@ -4342,10 +4342,10 @@ static void emitTargetTeamsDistributePar
// Emit teams region as a standalone region.
auto &&CodeGenTeams = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
PrePostActionTy &Action) {
+ Action.Enter(CGF);
CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
CGF.CGM.getOpenMPRuntime().emitInlinedDirective(
CGF, OMPD_distribute, CodeGenDistribute, /*HasCancel=*/false);
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
@@ -4614,12 +4614,12 @@ static void emitTargetParallelRegion(Cod
auto *CS = S.getCapturedStmt(OMPD_parallel);
Action.Enter(CGF);
auto &&CodeGen = [&S, CS](CodeGenFunction &CGF, PrePostActionTy &Action) {
+ Action.Enter(CGF);
CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
(void)CGF.EmitOMPFirstprivateClause(S, PrivateScope);
CGF.EmitOMPPrivateClause(S, PrivateScope);
CGF.EmitOMPReductionClauseInit(S, PrivateScope);
(void)PrivateScope.Privatize();
- Action.Enter(CGF);
// TODO: Add support for clauses.
CGF.EmitStmt(CS->getCapturedStmt());
CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_parallel);
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=327985&r1=327984&r2=327985&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Tue Mar 20 07:45:59 2018
@@ -803,7 +803,8 @@ void DSAStackTy::addDSA(ValueDecl *D, Ex
/// \brief Build a variable declaration for OpenMP loop iteration variable.
static VarDecl *buildVarDecl(Sema &SemaRef, SourceLocation Loc, QualType Type,
- StringRef Name, const AttrVec *Attrs = nullptr) {
+ StringRef Name, const AttrVec *Attrs = nullptr,
+ DeclRefExpr *OrigRef = nullptr) {
DeclContext *DC = SemaRef.CurContext;
IdentifierInfo *II = &SemaRef.PP.getIdentifierTable().get(Name);
TypeSourceInfo *TInfo = SemaRef.Context.getTrivialTypeSourceInfo(Type, Loc);
@@ -815,6 +816,10 @@ static VarDecl *buildVarDecl(Sema &SemaR
Decl->addAttr(*I);
}
Decl->setImplicit();
+ if (OrigRef) {
+ Decl->addAttr(
+ OMPReferencedVarAttr::CreateImplicit(SemaRef.Context, OrigRef));
+ }
return Decl;
}
@@ -1462,7 +1467,11 @@ void Sema::setOpenMPCaptureKind(FieldDec
}
if (DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective,
NewLevel)) {
- OMPC = OMPC_firstprivate;
+ OMPC = OMPC_map;
+ if (D->getType()->isScalarType() &&
+ DSAStack->getDefaultDMAAtLevel(NewLevel) !=
+ DefaultMapAttributes::DMA_tofrom_scalar)
+ OMPC = OMPC_firstprivate;
break;
}
}
@@ -1525,7 +1534,7 @@ void Sema::EndOpenMPDSABlock(Stmt *CurDi
// region uses original variable for proper diagnostics.
auto *VDPrivate = buildVarDecl(
*this, DE->getExprLoc(), Type.getUnqualifiedType(),
- VD->getName(), VD->hasAttrs() ? &VD->getAttrs() : nullptr);
+ VD->getName(), VD->hasAttrs() ? &VD->getAttrs() : nullptr, DRE);
ActOnUninitializedDecl(VDPrivate);
if (VDPrivate->isInvalidDecl())
continue;
@@ -4206,9 +4215,12 @@ DeclRefExpr *OpenMPIterationSpaceChecker
Expr *OpenMPIterationSpaceChecker::BuildPrivateCounterVar() const {
if (LCDecl && !LCDecl->isInvalidDecl()) {
auto Type = LCDecl->getType().getNonReferenceType();
- auto *PrivateVar =
- buildVarDecl(SemaRef, DefaultLoc, Type, LCDecl->getName(),
- LCDecl->hasAttrs() ? &LCDecl->getAttrs() : nullptr);
+ auto *PrivateVar = buildVarDecl(
+ SemaRef, DefaultLoc, Type, LCDecl->getName(),
+ LCDecl->hasAttrs() ? &LCDecl->getAttrs() : nullptr,
+ isa<VarDecl>(LCDecl)
+ ? buildDeclRefExpr(SemaRef, cast<VarDecl>(LCDecl), Type, DefaultLoc)
+ : nullptr);
if (PrivateVar->isInvalidDecl())
return nullptr;
return buildDeclRefExpr(SemaRef, PrivateVar, Type, DefaultLoc);
@@ -9322,8 +9334,10 @@ OMPClause *Sema::ActOnOpenMPPrivateClaus
// IdResolver, so the code in the OpenMP region uses original variable for
// proper diagnostics.
Type = Type.getUnqualifiedType();
- auto VDPrivate = buildVarDecl(*this, ELoc, Type, D->getName(),
- D->hasAttrs() ? &D->getAttrs() : nullptr);
+ auto VDPrivate =
+ buildVarDecl(*this, ELoc, Type, D->getName(),
+ D->hasAttrs() ? &D->getAttrs() : nullptr,
+ VD ? cast<DeclRefExpr>(SimpleRefExpr) : nullptr);
ActOnUninitializedDecl(VDPrivate);
if (VDPrivate->isInvalidDecl())
continue;
@@ -9561,8 +9575,10 @@ OMPClause *Sema::ActOnOpenMPFirstprivate
}
Type = Type.getUnqualifiedType();
- auto VDPrivate = buildVarDecl(*this, ELoc, Type, D->getName(),
- D->hasAttrs() ? &D->getAttrs() : nullptr);
+ auto VDPrivate =
+ buildVarDecl(*this, ELoc, Type, D->getName(),
+ D->hasAttrs() ? &D->getAttrs() : nullptr,
+ VD ? cast<DeclRefExpr>(SimpleRefExpr) : nullptr);
// Generate helper private variable and initialize it with the value of the
// original variable. The address of the original variable is replaced by
// the address of the new private variable in the CodeGen. This new variable
@@ -10454,8 +10470,10 @@ static bool ActOnOMPReductionKindClause(
Context.getAsArrayType(D->getType().getNonReferenceType()))
PrivateTy = D->getType().getNonReferenceType();
// Private copy.
- auto *PrivateVD = buildVarDecl(S, ELoc, PrivateTy, D->getName(),
- D->hasAttrs() ? &D->getAttrs() : nullptr);
+ auto *PrivateVD =
+ buildVarDecl(S, ELoc, PrivateTy, D->getName(),
+ D->hasAttrs() ? &D->getAttrs() : nullptr,
+ VD ? cast<DeclRefExpr>(SimpleRefExpr) : nullptr);
// Add initializer for private variable.
Expr *Init = nullptr;
auto *LHSDRE = buildDeclRefExpr(S, LHSVD, Type, ELoc);
@@ -10911,8 +10929,10 @@ OMPClause *Sema::ActOnOpenMPLinearClause
Type = Type.getNonReferenceType().getUnqualifiedType().getCanonicalType();
// Build private copy of original var.
- auto *Private = buildVarDecl(*this, ELoc, Type, D->getName(),
- D->hasAttrs() ? &D->getAttrs() : nullptr);
+ auto *Private =
+ buildVarDecl(*this, ELoc, Type, D->getName(),
+ D->hasAttrs() ? &D->getAttrs() : nullptr,
+ VD ? cast<DeclRefExpr>(SimpleRefExpr) : nullptr);
auto *PrivateRef = buildDeclRefExpr(*this, Private, Type, ELoc);
// Build var to save initial value.
VarDecl *Init = buildVarDecl(*this, ELoc, Type, ".linear.start");
@@ -13072,8 +13092,10 @@ OMPClause *Sema::ActOnOpenMPUseDevicePtr
}
// Build the private variable and the expression that refers to it.
- auto VDPrivate = buildVarDecl(*this, ELoc, Type, D->getName(),
- D->hasAttrs() ? &D->getAttrs() : nullptr);
+ auto VDPrivate =
+ buildVarDecl(*this, ELoc, Type, D->getName(),
+ D->hasAttrs() ? &D->getAttrs() : nullptr,
+ VD ? cast<DeclRefExpr>(SimpleRefExpr) : nullptr);
if (VDPrivate->isInvalidDecl())
continue;
Modified: cfe/trunk/test/OpenMP/nvptx_data_sharing.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_data_sharing.cpp?rev=327985&r1=327984&r2=327985&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_data_sharing.cpp (original)
+++ cfe/trunk/test/OpenMP/nvptx_data_sharing.cpp Tue Mar 20 07:45:59 2018
@@ -18,8 +18,10 @@ void test_ds(){
a = 1000;
}
int b = 100;
- #pragma omp parallel
+ int c = 1000;
+ #pragma omp parallel private(c)
{
+ int *c1 = &c;
b = a + 10000;
}
}
@@ -73,6 +75,15 @@ void test_ds(){
// CK1: [[SHARGSTMP16:%.+]] = load i32*, i32** [[SHARGSTMP15]]
// CK1: call void @__omp_outlined__{{.*}}({{.*}}, i32* [[SHARGSTMP16]])
+/// outlined function for the second parallel region ///
+
+// CK1: define internal void @{{.+}}(i32* noalias %{{.+}}, i32* noalias %{{.+}}, i32* dereferenceable{{.+}}, i32* dereferenceable{{.+}})
+// CK1: [[RES:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i64 4, i16 0)
+// CK1: [[GLOBALS:%.+]] = bitcast i8* [[RES]] to [[GLOBAL_TY:%.+]]*
+// CK1: [[C_ADDR:%.+]] = getelementptr inbounds [[GLOBAL_TY]], [[GLOBAL_TY]]* [[GLOBALS]], i32 0, i32 0
+// CK1: store i32* [[C_ADDR]], i32** %
+// CK1: call void @__kmpc_data_sharing_pop_stack(i8* [[RES]])
+
/// ========= In the data sharing wrapper function ========= ///
// CK1: {{.*}}define internal void @__omp_outlined{{.*}}wrapper({{.*}})
More information about the cfe-commits
mailing list