From 08aa612e2647cb62f5b6a6123aef4a1f28f72b00 Mon Sep 17 00:00:00 2001
From: Alexey Bataev <a.bataev@hotmail.com>
Date: Wed, 2 May 2018 14:20:50 +0000
Subject: [PATCH] [OPENMP] Emit names of the globals depending on target.
Some symbols are not allowed to be used as names on some targets. Patch
ries to unify the emission of the names of LLVM globals so they could be
used on different targets.
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@331358 91177308-0d34-0410-b5e6-96231b3b80d8
---
lib/CodeGen/CGOpenMPRuntime.cpp | 190 +++++++++++++++----------
lib/CodeGen/CGOpenMPRuntime.h | 11 +-
lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp | 3 +-
test/OpenMP/nvptx_parallel_codegen.cpp | 12 +-
4 files changed, 141 insertions(+), 75 deletions(-)
diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp
index b7e4cb46c8..2c66f6947d 100644
--- a/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -783,9 +783,10 @@ static void emitInitWithReductionInitializer(CodeGenFunction &CGF,
CGF.EmitIgnoredExpr(InitOp);
} else {
llvm::Constant *Init = CGF.CGM.EmitNullConstant(Ty);
+ std::string Name = CGF.CGM.getOpenMPRuntime().getName({"init"});
auto *GV = new llvm::GlobalVariable(
CGF.CGM.getModule(), Init->getType(), /*isConstant=*/true,
- llvm::GlobalValue::PrivateLinkage, Init, ".init");
+ llvm::GlobalValue::PrivateLinkage, Init, Name);
LValue LV = CGF.MakeNaturalAlignAddrLValue(GV, Ty);
RValue InitRVal;
switch (CGF.getEvaluationKind(Ty)) {
@@ -1216,8 +1217,10 @@ static FieldDecl *addFieldToRecordDecl(ASTContext &C, DeclContext *DC,
return Field;
}
-CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM)
- : CGM(CGM), OffloadEntriesInfoManager(CGM) {
+CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM, StringRef FirstSeparator,
+ StringRef Separator)
+ : CGM(CGM), FirstSeparator(FirstSeparator), Separator(Separator),
+ OffloadEntriesInfoManager(CGM) {
ASTContext &C = CGM.getContext();
RecordDecl *RD = C.buildImplicitRecord("ident_t");
QualType KmpInt32Ty = C.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/1);
@@ -1244,6 +1247,17 @@ void CGOpenMPRuntime::clear() {
InternalVars.clear();
}
+std::string CGOpenMPRuntime::getName(ArrayRef<StringRef> Parts) const {
+ SmallString<128> Buffer;
+ llvm::raw_svector_ostream OS(Buffer);
+ StringRef Sep = FirstSeparator;
+ for (StringRef Part : Parts) {
+ OS << Sep << Part;
+ Sep = Separator;
+ }
+ return OS.str();
+}
+
static llvm::Function *
emitCombinerOrInitializer(CodeGenModule &CGM, QualType Ty,
const Expr *CombinerInitializer, const VarDecl *In,
@@ -1261,9 +1275,10 @@ emitCombinerOrInitializer(CodeGenModule &CGM, QualType Ty,
const CGFunctionInfo &FnInfo =
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
llvm::FunctionType *FnTy = CGM.getTypes().GetFunctionType(FnInfo);
- auto *Fn = llvm::Function::Create(
- FnTy, llvm::GlobalValue::InternalLinkage,
- IsCombiner ? ".omp_combiner." : ".omp_initializer.", &CGM.getModule());
+ std::string Name = CGM.getOpenMPRuntime().getName(
+ {IsCombiner ? "omp_combiner" : "omp_initializer", ""});
+ auto *Fn = llvm::Function::Create(FnTy, llvm::GlobalValue::InternalLinkage,
+ Name, &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, FnInfo);
Fn->removeFnAttr(llvm::Attribute::NoInline);
Fn->removeFnAttr(llvm::Attribute::OptimizeNone);
@@ -2434,8 +2449,9 @@ CGOpenMPRuntime::getOrCreateThreadPrivateCache(const VarDecl *VD) {
assert(!CGM.getLangOpts().OpenMPUseTLS ||
!CGM.getContext().getTargetInfo().isTLSSupported());
// Lookup the entry, lazily creating it if necessary.
- return getOrCreateInternalVariable(CGM.Int8PtrPtrTy,
- Twine(CGM.getMangledName(VD), ".cache."));
+ std::string Suffix = getName({"cache", ""});
+ return getOrCreateInternalVariable(
+ CGM.Int8PtrPtrTy, Twine(CGM.getMangledName(VD)).concat(Suffix));
}
Address CGOpenMPRuntime::getAddrOfThreadPrivate(CodeGenFunction &CGF,
@@ -2501,8 +2517,9 @@ llvm::Function *CGOpenMPRuntime::emitThreadPrivateVarDefinition(
const auto &FI = CGM.getTypes().arrangeBuiltinFunctionDeclaration(
CGM.getContext().VoidPtrTy, Args);
llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
- llvm::Function *Fn = CGM.CreateGlobalInitOrDestructFunction(
- FTy, ".__kmpc_global_ctor_.", FI, Loc);
+ std::string Name = getName({"__kmpc_global_ctor_", ""});
+ llvm::Function *Fn =
+ CGM.CreateGlobalInitOrDestructFunction(FTy, Name, FI, Loc);
CtorCGF.StartFunction(GlobalDecl(), CGM.getContext().VoidPtrTy, Fn, FI,
Args, Loc, Loc);
llvm::Value *ArgVal = CtorCGF.EmitLoadOfScalar(
@@ -2533,8 +2550,9 @@ llvm::Function *CGOpenMPRuntime::emitThreadPrivateVarDefinition(
const auto &FI = CGM.getTypes().arrangeBuiltinFunctionDeclaration(
CGM.getContext().VoidTy, Args);
llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
- llvm::Function *Fn = CGM.CreateGlobalInitOrDestructFunction(
- FTy, ".__kmpc_global_dtor_.", FI, Loc);
+ std::string Name = getName({"__kmpc_global_dtor_", ""});
+ llvm::Function *Fn =
+ CGM.CreateGlobalInitOrDestructFunction(FTy, Name, FI, Loc);
auto NL = ApplyDebugLocation::CreateEmpty(DtorCGF);
DtorCGF.StartFunction(GlobalDecl(), CGM.getContext().VoidTy, Fn, FI, Args,
Loc, Loc);
@@ -2576,9 +2594,9 @@ llvm::Function *CGOpenMPRuntime::emitThreadPrivateVarDefinition(
if (!CGF) {
auto *InitFunctionTy =
llvm::FunctionType::get(CGM.VoidTy, /*isVarArg*/ false);
+ std::string Name = getName({"__omp_threadprivate_init_", ""});
llvm::Function *InitFunction = CGM.CreateGlobalInitOrDestructFunction(
- InitFunctionTy, ".__omp_threadprivate_init_.",
- CGM.getTypes().arrangeNullaryFunction());
+ InitFunctionTy, Name, CGM.getTypes().arrangeNullaryFunction());
CodeGenFunction InitCGF(CGM);
FunctionArgList ArgList;
InitCGF.StartFunction(GlobalDecl(), CGM.getContext().VoidTy, InitFunction,
@@ -2728,16 +2746,19 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
Address CGOpenMPRuntime::getAddrOfArtificialThreadPrivate(CodeGenFunction &CGF,
QualType VarType,
StringRef Name) {
- llvm::Twine VarName(Name, ".artificial.");
+ std::string Suffix = getName({"artificial", ""});
+ std::string CacheSuffix = getName({"cache", ""});
llvm::Type *VarLVType = CGF.ConvertTypeForMem(VarType);
- llvm::Value *GAddr = getOrCreateInternalVariable(VarLVType, VarName);
+ llvm::Value *GAddr =
+ getOrCreateInternalVariable(VarLVType, Twine(Name).concat(Suffix));
llvm::Value *Args[] = {
emitUpdateLocation(CGF, SourceLocation()),
getThreadID(CGF, SourceLocation()),
CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(GAddr, CGM.VoidPtrTy),
CGF.Builder.CreateIntCast(CGF.getTypeSize(VarType), CGM.SizeTy,
/*IsSigned=*/false),
- getOrCreateInternalVariable(CGM.VoidPtrPtrTy, VarName + ".cache.")};
+ getOrCreateInternalVariable(
+ CGM.VoidPtrPtrTy, Twine(Name).concat(Suffix).concat(CacheSuffix))};
return Address(
CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
CGF.EmitRuntimeCall(
@@ -2826,9 +2847,8 @@ void CGOpenMPRuntime::emitParallelCall(CodeGenFunction &CGF, SourceLocation Loc,
// OutlinedFn(>id, &zero, CapturedStruct);
Address ThreadIDAddr = RT.emitThreadIDAddress(CGF, Loc);
- Address ZeroAddr =
- CGF.CreateTempAlloca(CGF.Int32Ty, CharUnits::fromQuantity(4),
- /*Name*/ ".zero.addr");
+ Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
+ /*Name*/ ".zero.addr");
CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0));
llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
OutlinedFnArgs.push_back(ThreadIDAddr.getPointer());
@@ -2894,8 +2914,9 @@ CGOpenMPRuntime::getOrCreateInternalVariable(llvm::Type *Ty,
}
llvm::Value *CGOpenMPRuntime::getCriticalRegionLock(StringRef CriticalName) {
- llvm::Twine Name(".gomp_critical_user_", CriticalName);
- return getOrCreateInternalVariable(KmpCriticalNameTy, Name.concat(".var"));
+ std::string Prefix = Twine("gomp_critical_user_", CriticalName).str();
+ std::string Name = getName({Prefix, "var"});
+ return getOrCreateInternalVariable(KmpCriticalNameTy, Name);
}
namespace {
@@ -3042,9 +3063,11 @@ static llvm::Value *emitCopyprivateCopyFunction(
Args.push_back(&RHSArg);
const auto &CGFI =
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
- auto *Fn = llvm::Function::Create(
- CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
- ".omp.copyprivate.copy_func", &CGM.getModule());
+ std::string Name =
+ CGM.getOpenMPRuntime().getName({"omp", "copyprivate", "copy_func"});
+ auto *Fn = llvm::Function::Create(CGM.getTypes().GetFunctionType(CGFI),
+ llvm::GlobalValue::InternalLinkage, Name,
+ &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
Fn->setDoesNotRecurse();
CodeGenFunction CGF(CGM);
@@ -3712,14 +3735,16 @@ CGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
// host entries section. These will be defined by the linker.
llvm::Type *OffloadEntryTy =
CGM.getTypes().ConvertTypeForMem(getTgtOffloadEntryQTy());
+ std::string EntriesBeginName = getName({"omp_offloading", "entries_begin"});
auto *HostEntriesBegin = new llvm::GlobalVariable(
M, OffloadEntryTy, /*isConstant=*/true,
llvm::GlobalValue::ExternalLinkage, /*Initializer=*/nullptr,
- ".omp_offloading.entries_begin");
- auto *HostEntriesEnd = new llvm::GlobalVariable(
- M, OffloadEntryTy, /*isConstant=*/true,
- llvm::GlobalValue::ExternalLinkage, /*Initializer=*/nullptr,
- ".omp_offloading.entries_end");
+ EntriesBeginName);
+ std::string EntriesEndName = getName({"omp_offloading", "entries_end"});
+ auto *HostEntriesEnd =
+ new llvm::GlobalVariable(M, OffloadEntryTy, /*isConstant=*/true,
+ llvm::GlobalValue::ExternalLinkage,
+ /*Initializer=*/nullptr, EntriesEndName);
// Create all device images
auto *DeviceImageTy = cast<llvm::StructType>(
@@ -3730,12 +3755,14 @@ CGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
for (const llvm::Triple &Device : Devices) {
StringRef T = Device.getTriple();
+ std::string BeginName = getName({"omp_offloading", "img_start", ""});
auto *ImgBegin = new llvm::GlobalVariable(
M, CGM.Int8Ty, /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage,
- /*Initializer=*/nullptr, Twine(".omp_offloading.img_start.", T));
+ /*Initializer=*/nullptr, Twine(BeginName).concat(T));
+ std::string EndName = getName({"omp_offloading", "img_end", ""});
auto *ImgEnd = new llvm::GlobalVariable(
M, CGM.Int8Ty, /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage,
- /*Initializer=*/nullptr, Twine(".omp_offloading.img_end.", T));
+ /*Initializer=*/nullptr, Twine(EndName).concat(T));
llvm::Constant *Data[] = {ImgBegin, ImgEnd, HostEntriesBegin,
HostEntriesEnd};
@@ -3744,10 +3771,11 @@ CGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
}
// Create device images global array.
+ std::string ImagesName = getName({"omp_offloading", "device_images"});
llvm::GlobalVariable *DeviceImages =
- DeviceImagesEntries.finishAndCreateGlobal(".omp_offloading.device_images",
- CGM.getPointerAlign(),
- /*isConstant=*/true);
+ DeviceImagesEntries.finishAndCreateGlobal(ImagesName,
+ CGM.getPointerAlign(),
+ /*isConstant=*/true);
DeviceImages->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
// This is a Zero array to be used in the creation of the constant expressions
@@ -3760,8 +3788,9 @@ CGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
llvm::ConstantExpr::getGetElementPtr(DeviceImages->getValueType(),
DeviceImages, Index),
HostEntriesBegin, HostEntriesEnd};
+ std::string Descriptor = getName({"omp_offloading", "descriptor"});
llvm::GlobalVariable *Desc = createConstantGlobalStruct(
- CGM, getTgtBinaryDescriptorQTy(), Data, ".omp_offloading.descriptor");
+ CGM, getTgtBinaryDescriptorQTy(), Data, Descriptor);
// Emit code to register or unregister the descriptor at execution
// startup or closing, respectively.
@@ -3779,8 +3808,8 @@ CGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
const auto &FI =
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
- UnRegFn = CGM.CreateGlobalInitOrDestructFunction(
- FTy, ".omp_offloading.descriptor_unreg", FI);
+ std::string UnregName = getName({"omp_offloading", "descriptor_unreg"});
+ UnRegFn = CGM.CreateGlobalInitOrDestructFunction(FTy, UnregName, FI);
CGF.StartFunction(GlobalDecl(), C.VoidTy, UnRegFn, FI, Args);
CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_unregister_lib),
Desc);
@@ -3794,8 +3823,8 @@ CGOpenMPRuntime::createOffloadingBinaryDescriptorRegistration() {
CGF.disableDebugInfo();
const auto &FI = CGM.getTypes().arrangeNullaryFunction();
llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
- RegFn = CGM.CreateGlobalInitOrDestructFunction(
- FTy, ".omp_offloading.descriptor_reg", FI);
+ std::string Descriptor = getName({"omp_offloading", "descriptor_reg"});
+ RegFn = CGM.CreateGlobalInitOrDestructFunction(FTy, Descriptor, FI);
CGF.StartFunction(GlobalDecl(), C.VoidTy, RegFn, FI, FunctionArgList());
CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_register_lib), Desc);
// Create a variable to drive the registration and unregistration of the
@@ -3832,10 +3861,10 @@ void CGOpenMPRuntime::createOffloadEntry(
// Create constant string with the name.
llvm::Constant *StrPtrInit = llvm::ConstantDataArray::getString(C, Name);
- auto *Str =
- new llvm::GlobalVariable(M, StrPtrInit->getType(), /*isConstant=*/true,
- llvm::GlobalValue::InternalLinkage, StrPtrInit,
- ".omp_offloading.entry_name");
+ std::string StringName = getName({"omp_offloading", "entry_name"});
+ auto *Str = new llvm::GlobalVariable(
+ M, StrPtrInit->getType(), /*isConstant=*/true,
+ llvm::GlobalValue::InternalLinkage, StrPtrInit, StringName);
Str->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
llvm::Constant *Data[] = {llvm::ConstantExpr::getBitCast(ID, CGM.VoidPtrTy),
@@ -3843,12 +3872,14 @@ void CGOpenMPRuntime::createOffloadEntry(
llvm::ConstantInt::get(CGM.SizeTy, Size),
llvm::ConstantInt::get(CGM.Int32Ty, Flags),
llvm::ConstantInt::get(CGM.Int32Ty, 0)};
- llvm::GlobalVariable *Entry = createConstantGlobalStruct(
- CGM, getTgtOffloadEntryQTy(), Data, Twine(".omp_offloading.entry.", Name),
- Linkage);
+ std::string EntryName = getName({"omp_offloading", "entry", ""});
+ llvm::GlobalVariable *Entry =
+ createConstantGlobalStruct(CGM, getTgtOffloadEntryQTy(), Data,
+ Twine(EntryName).concat(Name), Linkage);
// The entry has to be created in the section the linker expects it to be.
- Entry->setSection(".omp_offloading.entries");
+ std::string Section = getName({"omp_offloading", "entries"});
+ Entry->setSection(Section);
}
void CGOpenMPRuntime::createOffloadEntriesAndInfoMetadata() {
@@ -4267,9 +4298,9 @@ emitProxyTaskFunction(CodeGenModule &CGM, SourceLocation Loc,
CGM.getTypes().arrangeBuiltinFunctionDeclaration(KmpInt32Ty, Args);
llvm::FunctionType *TaskEntryTy =
CGM.getTypes().GetFunctionType(TaskEntryFnInfo);
- auto *TaskEntry =
- llvm::Function::Create(TaskEntryTy, llvm::GlobalValue::InternalLinkage,
- ".omp_task_entry.", &CGM.getModule());
+ std::string Name = CGM.getOpenMPRuntime().getName({"omp_task_entry", ""});
+ auto *TaskEntry = llvm::Function::Create(
+ TaskEntryTy, llvm::GlobalValue::InternalLinkage, Name, &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), TaskEntry, TaskEntryFnInfo);
TaskEntry->setDoesNotRecurse();
CodeGenFunction CGF(CGM);
@@ -4369,9 +4400,11 @@ static llvm::Value *emitDestructorsFunction(CodeGenModule &CGM,
CGM.getTypes().arrangeBuiltinFunctionDeclaration(KmpInt32Ty, Args);
llvm::FunctionType *DestructorFnTy =
CGM.getTypes().GetFunctionType(DestructorFnInfo);
+ std::string Name =
+ CGM.getOpenMPRuntime().getName({"omp_task_destructor", ""});
auto *DestructorFn =
llvm::Function::Create(DestructorFnTy, llvm::GlobalValue::InternalLinkage,
- ".omp_task_destructor.", &CGM.getModule());
+ Name, &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), DestructorFn,
DestructorFnInfo);
DestructorFn->setDoesNotRecurse();
@@ -4461,9 +4494,11 @@ emitTaskPrivateMappingFunction(CodeGenModule &CGM, SourceLocation Loc,
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
llvm::FunctionType *TaskPrivatesMapTy =
CGM.getTypes().GetFunctionType(TaskPrivatesMapFnInfo);
+ std::string Name =
+ CGM.getOpenMPRuntime().getName({"omp_task_privates_map", ""});
auto *TaskPrivatesMap = llvm::Function::Create(
- TaskPrivatesMapTy, llvm::GlobalValue::InternalLinkage,
- ".omp_task_privates_map.", &CGM.getModule());
+ TaskPrivatesMapTy, llvm::GlobalValue::InternalLinkage, Name,
+ &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), TaskPrivatesMap,
TaskPrivatesMapFnInfo);
TaskPrivatesMap->removeFnAttr(llvm::Attribute::NoInline);
@@ -4653,9 +4688,9 @@ emitTaskDupFunction(CodeGenModule &CGM, SourceLocation Loc,
const auto &TaskDupFnInfo =
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
llvm::FunctionType *TaskDupTy = CGM.getTypes().GetFunctionType(TaskDupFnInfo);
- auto *TaskDup =
- llvm::Function::Create(TaskDupTy, llvm::GlobalValue::InternalLinkage,
- ".omp_task_dup.", &CGM.getModule());
+ std::string Name = CGM.getOpenMPRuntime().getName({"omp_task_dup", ""});
+ auto *TaskDup = llvm::Function::Create(
+ TaskDupTy, llvm::GlobalValue::InternalLinkage, Name, &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), TaskDup, TaskDupFnInfo);
TaskDup->setDoesNotRecurse();
CodeGenFunction CGF(CGM);
@@ -5306,9 +5341,10 @@ llvm::Value *CGOpenMPRuntime::emitReductionFunction(
Args.push_back(&RHSArg);
const auto &CGFI =
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
- auto *Fn = llvm::Function::Create(
- CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
- ".omp.reduction.reduction_func", &CGM.getModule());
+ std::string Name = getName({"omp", "reduction", "reduction_func"});
+ auto *Fn = llvm::Function::Create(CGM.getTypes().GetFunctionType(CGFI),
+ llvm::GlobalValue::InternalLinkage, Name,
+ &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
Fn->setDoesNotRecurse();
CodeGenFunction CGF(CGM);
@@ -5510,7 +5546,8 @@ void CGOpenMPRuntime::emitReduction(CodeGenFunction &CGF, SourceLocation Loc,
Privates, LHSExprs, RHSExprs, ReductionOps);
// 3. Create static kmp_critical_name lock = { 0 };
- llvm::Value *Lock = getCriticalRegionLock(".reduction");
+ std::string Name = getName({"reduction"});
+ llvm::Value *Lock = getCriticalRegionLock(Name);
// 4. Build res = __kmpc_reduce{_nowait}(<loc>, <gtid>, <n>, sizeof(RedList),
// RedList, reduce_func, &<lock>);
@@ -5659,10 +5696,11 @@ void CGOpenMPRuntime::emitReduction(CodeGenFunction &CGF, SourceLocation Loc,
} else {
// Emit as a critical region.
auto &&CritRedGen = [E, Loc](CodeGenFunction &CGF, const Expr *,
- const Expr *, const Expr *) {
+ const Expr *, const Expr *) {
CGOpenMPRuntime &RT = CGF.CGM.getOpenMPRuntime();
+ std::string Name = RT.getName({"atomic_reduction"});
RT.emitCriticalRegion(
- CGF, ".atomic_reduction",
+ CGF, Name,
[=](CodeGenFunction &CGF, PrePostActionTy &Action) {
Action.Enter(CGF);
emitReductionCombiner(CGF, E);
@@ -5717,9 +5755,10 @@ static std::string generateUniqueName(CodeGenModule &CGM, StringRef Prefix,
if (!D)
D = cast<VarDecl>(cast<DeclRefExpr>(Ref)->getDecl());
D = D->getCanonicalDecl();
- Out << Prefix << "."
- << (D->isLocalVarDeclOrParm() ? D->getName() : CGM.getMangledName(D))
- << "_" << D->getCanonicalDecl()->getLocStart().getRawEncoding();
+ std::string Name = CGM.getOpenMPRuntime().getName(
+ {D->isLocalVarDeclOrParm() ? D->getName() : CGM.getMangledName(D)});
+ Out << Prefix << Name << "_"
+ << D->getCanonicalDecl()->getLocStart().getRawEncoding();
return Out.str();
}
@@ -5742,8 +5781,9 @@ static llvm::Value *emitReduceInitFunction(CodeGenModule &CGM,
const auto &FnInfo =
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
llvm::FunctionType *FnTy = CGM.getTypes().GetFunctionType(FnInfo);
+ std::string Name = CGM.getOpenMPRuntime().getName({"red_init", ""});
auto *Fn = llvm::Function::Create(FnTy, llvm::GlobalValue::InternalLinkage,
- ".red_init.", &CGM.getModule());
+ Name, &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, FnInfo);
Fn->setDoesNotRecurse();
CodeGenFunction CGF(CGM);
@@ -5818,8 +5858,9 @@ static llvm::Value *emitReduceCombFunction(CodeGenModule &CGM,
const auto &FnInfo =
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
llvm::FunctionType *FnTy = CGM.getTypes().GetFunctionType(FnInfo);
+ std::string Name = CGM.getOpenMPRuntime().getName({"red_comb", ""});
auto *Fn = llvm::Function::Create(FnTy, llvm::GlobalValue::InternalLinkage,
- ".red_comb.", &CGM.getModule());
+ Name, &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, FnInfo);
Fn->setDoesNotRecurse();
CodeGenFunction CGF(CGM);
@@ -5887,8 +5928,9 @@ static llvm::Value *emitReduceFiniFunction(CodeGenModule &CGM,
const auto &FnInfo =
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
llvm::FunctionType *FnTy = CGM.getTypes().GetFunctionType(FnInfo);
+ std::string Name = CGM.getOpenMPRuntime().getName({"red_fini", ""});
auto *Fn = llvm::Function::Create(FnTy, llvm::GlobalValue::InternalLinkage,
- ".red_fini.", &CGM.getModule());
+ Name, &CGM.getModule());
CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, FnInfo);
Fn->setDoesNotRecurse();
CodeGenFunction CGF(CGM);
@@ -6252,10 +6294,11 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
OutlinedFn->setLinkage(llvm::GlobalValue::ExternalLinkage);
OutlinedFn->setDSOLocal(false);
} else {
+ std::string Name = getName({"omp_offload", "region_id"});
OutlinedFnID = new llvm::GlobalVariable(
CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
llvm::GlobalValue::PrivateLinkage,
- llvm::Constant::getNullValue(CGM.Int8Ty), ".omp_offload.region_id");
+ llvm::Constant::getNullValue(CGM.Int8Ty), Name);
}
// Register the information for the entry associated with this target region.
@@ -7292,10 +7335,11 @@ emitOffloadingArrays(CodeGenFunction &CGF,
auto *SizesArrayInit = llvm::ConstantArray::get(
llvm::ArrayType::get(CGM.SizeTy, ConstSizes.size()), ConstSizes);
+ std::string Name = CGM.getOpenMPRuntime().getName({"offload_sizes"});
auto *SizesArrayGbl = new llvm::GlobalVariable(
CGM.getModule(), SizesArrayInit->getType(),
/*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
- SizesArrayInit, ".offload_sizes");
+ SizesArrayInit, Name);
SizesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
Info.SizesArray = SizesArrayGbl;
}
@@ -7304,10 +7348,12 @@ emitOffloadingArrays(CodeGenFunction &CGF,
// fill arrays. Instead, we create an array constant.
llvm::Constant *MapTypesArrayInit =
llvm::ConstantDataArray::get(CGF.Builder.getContext(), MapTypes);
+ std::string MaptypesName =
+ CGM.getOpenMPRuntime().getName({"offload_maptypes"});
auto *MapTypesArrayGbl = new llvm::GlobalVariable(
CGM.getModule(), MapTypesArrayInit->getType(),
/*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
- MapTypesArrayInit, ".offload_maptypes");
+ MapTypesArrayInit, MaptypesName);
MapTypesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
Info.MapTypesArray = MapTypesArrayGbl;
diff --git a/lib/CodeGen/CGOpenMPRuntime.h b/lib/CodeGen/CGOpenMPRuntime.h
index 0dea165c13..e1ce09adc2 100644
--- a/lib/CodeGen/CGOpenMPRuntime.h
+++ b/lib/CodeGen/CGOpenMPRuntime.h
@@ -213,6 +213,11 @@ public:
protected:
CodeGenModule &CGM;
+ StringRef FirstSeparator, Separator;
+
+ /// Constructor allowing to redefine the name separator for the variables.
+ explicit CGOpenMPRuntime(CodeGenModule &CGM, StringRef FirstSeparator,
+ StringRef Separator);
/// \brief Creates offloading entry for the provided entry ID \a ID,
/// address \a Addr, size \a Size, and flags \a Flags.
@@ -724,10 +729,14 @@ private:
Address Shareds, const OMPTaskDataTy &Data);
public:
- explicit CGOpenMPRuntime(CodeGenModule &CGM);
+ explicit CGOpenMPRuntime(CodeGenModule &CGM)
+ : CGOpenMPRuntime(CGM, ".", ".") {}
virtual ~CGOpenMPRuntime() {}
virtual void clear();
+ /// Get the platform-specific name separator.
+ std::string getName(ArrayRef<StringRef> Parts) const;
+
/// Emit code for the specified user defined reduction construct.
virtual void emitUserDefinedReduction(CodeGenFunction *CGF,
const OMPDeclareReductionDecl *D);
diff --git a/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp b/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
index 59858d3430..22e1e48a22 100644
--- a/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
+++ b/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
@@ -1184,7 +1184,8 @@ void CGOpenMPRuntimeNVPTX::emitTargetOutlinedFunction(
}
CGOpenMPRuntimeNVPTX::CGOpenMPRuntimeNVPTX(CodeGenModule &CGM)
- : CGOpenMPRuntime(CGM), CurrentExecutionMode(ExecutionMode::Unknown) {
+ : CGOpenMPRuntime(CGM, "_", "$"),
+ CurrentExecutionMode(ExecutionMode::Unknown) {
if (!CGM.getLangOpts().OpenMPIsDevice)
llvm_unreachable("OpenMP NVPTX can only handle device code.");
}
diff --git a/test/OpenMP/nvptx_parallel_codegen.cpp b/test/OpenMP/nvptx_parallel_codegen.cpp
index 932454a6f8..f8f91e8724 100644
--- a/test/OpenMP/nvptx_parallel_codegen.cpp
+++ b/test/OpenMP/nvptx_parallel_codegen.cpp
@@ -51,6 +51,14 @@ tx ftemplate(int n) {
b[2] += 1;
}
+ #pragma omp target
+ {
+ #pragma omp parallel
+ {
+ #pragma omp critical
+ ++a;
+ }
+ }
return a;
}
@@ -62,7 +70,9 @@ int bar(int n){
return a;
}
- // CHECK-NOT: define {{.*}}void {{@__omp_offloading_.+template.+l17}}_worker()
+// CHECK: @"_gomp_critical_user_$var" = common global [8 x i32] zeroinitializer
+
+// CHECK-NOT: define {{.*}}void {{@__omp_offloading_.+template.+l17}}_worker()
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l26}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
--
2.40.0