r331358 - [OPENMP] Emit names of the globals depending on target.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Wed May 2 07:20:50 PDT 2018
Author: abataev
Date: Wed May 2 07:20:50 2018
New Revision: 331358
URL: http://llvm.org/viewvc/llvm-project?rev=331358&view=rev
Log:
[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.
Modified:
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
cfe/trunk/test/OpenMP/nvptx_parallel_codegen.cpp
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=331358&r1=331357&r2=331358&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Wed May 2 07:20:50 2018
@@ -783,9 +783,10 @@ static void emitInitWithReductionInitial
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(A
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
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::getOrCreateThreadPrivat
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::emitThr
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::emitThr
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::emitThr
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::emitDeclareTargetV
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(C
// 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::getOrCreateInternalVari
}
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 *emitCopyprivateCopyF
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::createOffloadingBinaryD
// 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::createOffloadingBinaryD
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::createOffloadingBinaryD
}
// 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::createOffloadingBinaryD
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::createOffloadingBinaryD
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::createOffloadingBinaryD
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
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 *emitDestructorsFunct
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(CodeGenMo
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,
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::emitReduct
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(Code
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(Code
} 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(Co
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 *emitReduceInitFuncti
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 *emitReduceCombFuncti
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 *emitReduceFiniFuncti
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::emitTargetOutlined
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 &CG
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 &CG
// 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;
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h?rev=331358&r1=331357&r2=331358&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h Wed May 2 07:20:50 2018
@@ -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);
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp?rev=331358&r1=331357&r2=331358&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp Wed May 2 07:20:50 2018
@@ -1184,7 +1184,8 @@ void CGOpenMPRuntimeNVPTX::emitTargetOut
}
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.");
}
Modified: cfe/trunk/test/OpenMP/nvptx_parallel_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_parallel_codegen.cpp?rev=331358&r1=331357&r2=331358&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_parallel_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/nvptx_parallel_codegen.cpp Wed May 2 07:20:50 2018
@@ -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,
More information about the cfe-commits
mailing list