ssquare08 created this revision.
ssquare08 added a reviewer: jhuber6.
Herald added subscribers: mattd, asavonic, guansong, yaxunl.
Herald added a project: All.
ssquare08 requested review of this revision.
Herald added a reviewer: jdoerfert.
Herald added subscribers: cfe-commits, sstefan1.
Herald added a project: clang.
This is to support cases where static globals are marked declare
target. By default these file static globals are not externally
visible but in order for OpenMP runtime to access these symbols,
this changes here makes them externally visisble unless they
have "hidden" visibility attribute.
Making them externally visible, however, leads to symbol conflict
when two files have variables with the same name. Thus, these
symbols needs to be mangled on the device side of the compilation.
In order to do so, the host side mangles the symbol names and
passes that metadata information to the device side. It also uses
these mangled names if offload entry table so that the OPenMP
runtime can find these symbols during registration.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D129694
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/CodeGen/CGOpenMPRuntime.h
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/CodeGen/TargetInfo.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/declare_target_codegen.cpp
clang/test/OpenMP/declare_target_link_codegen.cpp
clang/test/OpenMP/declare_target_visibility_codegen.cpp
clang/test/OpenMP/nvptx_allocate_codegen.cpp
clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
clang/test/OpenMP/target_update_messages.cpp
Index: clang/test/OpenMP/target_update_messages.cpp
===================================================================
--- clang/test/OpenMP/target_update_messages.cpp
+++ clang/test/OpenMP/target_update_messages.cpp
@@ -14,13 +14,6 @@
argc = x; // expected-warning {{variable 'x' is uninitialized when used here}}
}
-static int y;
-#pragma omp declare target(y)
-
-void yyy() {
-#pragma omp target update to(y) // expected-error {{the host cannot update a declare target variable that is not externally visible.}}
-}
-
int __attribute__((visibility("hidden"))) z;
#pragma omp declare target(z)
Index: clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
===================================================================
--- clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
+++ clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
@@ -15,7 +15,7 @@
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
-// DEVICE-DAG: [[C_ADDR:.+]] = internal global i32 0,
+// DEVICE-DAG: [[C_ADDR:.+]] = global i32 0,
// DEVICE-DAG: [[CD_ADDR:@.+]] ={{ protected | }}global %struct.S zeroinitializer,
// HOST-DAG: @[[C_ADDR:.+]] = internal global i32 0,
// HOST-DAG: @[[CD_ADDR:.+]] ={{( protected | dso_local)?}} global %struct.S zeroinitializer,
@@ -72,6 +72,8 @@
// DEVICE-DAG: call void
// DEVICE-DAG: ret void
+// HOST-DAG: @.omp_offloading.entry_name = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_ENTRY_NAME:c__static__.+]]\00"
+// HOST-DAG: @.omp_offloading.entry.[[C_ENTRY_NAME]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* bitcast (i32* @[[C_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name, i32 0, i32 0), i64 4, i32 0, i32 0 }, section "omp_offloading_entries", align 1
// HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_ADDR]]\00"
// HOST-DAG: @.omp_offloading.entry.[[CD_ADDR]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* bitcast (%struct.S* @[[CD_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section "omp_offloading_entries", align 1
// HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_CTOR]]\00"
@@ -97,8 +99,8 @@
// HOST: [[C:%.*]] = load i32, i32* @[[C_ADDR]],
// HOST: store i32 [[C]], i32* %
-// HOST-DAG: !{i32 1, !"[[CD_ADDR]]", i32 0, i32 {{[0-9]+}}}
-// HOST-DAG: !{i32 1, !"[[C_ADDR]]", i32 0, i32 {{[0-9]+}}}
+// HOST-DAG: !{i32 1, !"[[CD_ADDR]]", i32 0, i32 {{[0-9]+}}, !"cd"}
+// HOST-DAG: !{i32 1, !"[[C_ENTRY_NAME]]", i32 0, i32 {{[0-9]+}}, !"c"}
// DEVICE: !nvvm.annotations
// DEVICE-DAG: !{void ()* [[C_CTOR]], !"kernel", i32 1}
Index: clang/test/OpenMP/nvptx_allocate_codegen.cpp
===================================================================
--- clang/test/OpenMP/nvptx_allocate_codegen.cpp
+++ clang/test/OpenMP/nvptx_allocate_codegen.cpp
@@ -89,7 +89,7 @@
// CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 4
// CHECK1-NEXT: [[B:%.*]] = alloca double, align 8
// CHECK1-NEXT: store i32 0, i32* [[RETVAL]], align 4
-// CHECK1-NEXT: store i32 2, i32* @_ZZ4mainE1a, align 4
+// CHECK1-NEXT: store i32 2, i32* @a1, align 4
// CHECK1-NEXT: store double 3.000000e+00, double* [[B]], align 8
// CHECK1-NEXT: [[CALL:%.*]] = call noundef i32 @_Z3fooIiET_v() #[[ATTR7:[0-9]+]]
// CHECK1-NEXT: ret i32 [[CALL]]
Index: clang/test/OpenMP/declare_target_visibility_codegen.cpp
===================================================================
--- clang/test/OpenMP/declare_target_visibility_codegen.cpp
+++ clang/test/OpenMP/declare_target_visibility_codegen.cpp
@@ -8,8 +8,8 @@
// HOST: @[[X:.+]] = internal global i32 0, align 4
// HOST: @y = hidden global i32 0
// HOST: @z = global i32 0
-// HOST-NOT: @.omp_offloading.entry.c
-// HOST-NOT: @.omp_offloading.entry.x
+// HOST: @.omp_offloading.entry.c__static__{{[0-9a-z]+_[0-9a-z]+}}
+// HOST: @.omp_offloading.entry.x__static__{{[0-9a-z]+_[0-9a-z]+}}
// HOST-NOT: @.omp_offloading.entry.y
// HOST: @.omp_offloading.entry.z
C() : x(0) {}
Index: clang/test/OpenMP/declare_target_link_codegen.cpp
===================================================================
--- clang/test/OpenMP/declare_target_link_codegen.cpp
+++ clang/test/OpenMP/declare_target_link_codegen.cpp
@@ -85,5 +85,5 @@
// HOST: [[C:%.*]] = load i32, i32* @c,
// HOST: store i32 [[C]], i32* %
-// CHECK: !{i32 1, !"c_decl_tgt_ref_ptr", i32 1, i32 {{[0-9]+}}}
+// CHECK: !{i32 1, !"c_decl_tgt_ref_ptr", i32 1, i32 {{[0-9]+}}, !"c"}
#endif // HEADER
Index: clang/test/OpenMP/declare_target_codegen.cpp
===================================================================
--- clang/test/OpenMP/declare_target_codegen.cpp
+++ clang/test/OpenMP/declare_target_codegen.cpp
@@ -43,7 +43,7 @@
// CHECK-DAG: @d ={{ protected | }}global i32 0,
// CHECK-DAG: @c = external global i32,
// CHECK-DAG: @globals ={{ protected | }}global %struct.S zeroinitializer,
-// CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
+// CHECK-DAG: [[STAT:@stat__static__.+]] = internal global %struct.S zeroinitializer,
// CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
// CHECK-DAG: @out_decl_target ={{ protected | }}global i32 0,
// CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)],
@@ -247,8 +247,8 @@
// CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
-// CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}}
-// CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}}
+// CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}, !"aaa"}
+// CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}, !"ccc"}
// CHECK-DAG: !{{{.+}}virtual_foo
#ifdef OMP5
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -12961,20 +12961,18 @@
return hasClauses(Clauses, K) || hasClauses(Clauses, ClauseTypes...);
}
-/// Check if the variables in the mapping clause are externally visible.
+/// Check if the variables in the mapping clause have hidden visibility attribute
static bool isClauseMappable(ArrayRef<OMPClause *> Clauses) {
for (const OMPClause *C : Clauses) {
if (auto *TC = dyn_cast<OMPToClause>(C))
return llvm::all_of(TC->all_decls(), [](ValueDecl *VD) {
return !VD || !VD->hasAttr<OMPDeclareTargetDeclAttr>() ||
- (VD->isExternallyVisible() &&
- VD->getVisibility() != HiddenVisibility);
+ (VD->getVisibility() != HiddenVisibility);
});
else if (auto *FC = dyn_cast<OMPFromClause>(C))
return llvm::all_of(FC->all_decls(), [](ValueDecl *VD) {
return !VD || !VD->hasAttr<OMPDeclareTargetDeclAttr>() ||
- (VD->isExternallyVisible() &&
- VD->getVisibility() != HiddenVisibility);
+ (VD->getVisibility() != HiddenVisibility);
});
}
Index: clang/lib/CodeGen/TargetInfo.cpp
===================================================================
--- clang/lib/CodeGen/TargetInfo.cpp
+++ clang/lib/CodeGen/TargetInfo.cpp
@@ -7289,8 +7289,20 @@
void NVPTXTargetCodeGenInfo::setTargetAttributes(
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
+ // Make sure any variable with OpenMP declare target is visible to runtime
+ // except for those with hidden visibility
+ if ( D && isa<VarDecl>(D) && D->hasAttr<OMPDeclareTargetDeclAttr>() &&
+ isa<llvm::GlobalVariable>(GV) ) {
+ if ( (GV->hasInternalLinkage() ||
+ GV->hasPrivateLinkage()) &&
+ !GV->hasHiddenVisibility() ) {
+ GV->setLinkage(llvm::GlobalVariable::ExternalLinkage);
+ GV->setDSOLocal(false);
+ }
+ }
if (GV->isDeclaration())
return;
+
const VarDecl *VD = dyn_cast_or_null<VarDecl>(D);
if (VD) {
if (M.getLangOpts().CUDA) {
@@ -9424,6 +9436,18 @@
void AMDGPUTargetCodeGenInfo::setTargetAttributes(
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
+ // Make sure any variable with OpenMP declare target is visible to runtime
+ // except for those with hidden visibility
+ if ( D && isa<VarDecl>(D) && D->hasAttr<OMPDeclareTargetDeclAttr>() &&
+ isa<llvm::GlobalVariable>(GV) ) {
+ if ( (GV->hasInternalLinkage() ||
+ GV->hasPrivateLinkage()) &&
+ !GV->hasHiddenVisibility() ) {
+ GV->setLinkage(llvm::GlobalVariable::ExternalLinkage);
+ GV->setDSOLocal(false);
+ }
+ }
+
if (requiresAMDGPUProtectedVisibility(D, GV)) {
GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
GV->setDSOLocal(true);
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -1501,6 +1501,23 @@
const auto *ND = cast<NamedDecl>(GD.getDecl());
std::string MangledName = getMangledNameImpl(*this, GD, ND);
+ if ( getLangOpts().OpenMPIsDevice ) {
+ if ( isa<VarDecl>(GD.getDecl()) ) {
+ const auto *VD = dyn_cast<VarDecl>(GD.getDecl());
+ llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
+ OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
+
+ if ( Res && (*Res == OMPDeclareTargetDeclAttr::MT_To) &&
+ !getOpenMPRuntime().hasRequiresUnifiedSharedMemory() &&
+ !VD->isExternallyVisible() ) {
+ StringRef HostMangledName =
+ getOpenMPRuntime().getHostMangledDeclareTargetGlobal(VD->getName());
+ if (!HostMangledName.empty())
+ MangledName = HostMangledName.str();
+ }
+ }
+ }
+
// Ensure either we have different ABIs between host and device compilations,
// says host compilation following MSVC ABI but device compilation follows
// Itanium C++ ABI or, if they follow the same ABI, kernel names after
@@ -6886,6 +6903,7 @@
SM.getDiagnostics().Report(diag::err_cannot_open_file)
<< PLoc.getFilename() << EC.message();
}
+
OS << llvm::format("%x", ID.getFile()) << llvm::format("%x", ID.getDevice())
<< "_" << llvm::utohexstr(Result.low(), /*LowerCase=*/true, /*Width=*/8);
} else {
Index: clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
+++ clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
@@ -70,7 +70,8 @@
/// address \a Addr, size \a Size, and flags \a Flags.
void createOffloadEntry(llvm::Constant *ID, llvm::Constant *Addr,
uint64_t Size, int32_t Flags,
- llvm::GlobalValue::LinkageTypes Linkage) override;
+ llvm::GlobalValue::LinkageTypes Linkage,
+ StringRef MangledName) override;
/// Emit outlined function specialized for the Fork-Join
/// programming model for applicable target directives on the NVPTX device.
Index: clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -1122,7 +1122,8 @@
void CGOpenMPRuntimeGPU::createOffloadEntry(llvm::Constant *ID,
llvm::Constant *Addr,
uint64_t Size, int32_t,
- llvm::GlobalValue::LinkageTypes) {
+ llvm::GlobalValue::LinkageTypes,
+ StringRef) {
// TODO: Add support for global variables on the device after declare target
// support.
llvm::Function *Fn = dyn_cast<llvm::Function>(Addr);
Index: clang/lib/CodeGen/CGOpenMPRuntime.h
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.h
+++ clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -319,7 +319,8 @@
/// address \a Addr, size \a Size, and flags \a Flags.
virtual void createOffloadEntry(llvm::Constant *ID, llvm::Constant *Addr,
uint64_t Size, int32_t Flags,
- llvm::GlobalValue::LinkageTypes Linkage);
+ llvm::GlobalValue::LinkageTypes Linkage,
+ StringRef MangledName);
/// Helper to emit outlined function for 'target' directive.
/// \param D Directive to emit.
@@ -661,19 +662,22 @@
/// Type of the global variable.
CharUnits VarSize;
llvm::GlobalValue::LinkageTypes Linkage;
+ StringRef OrigName;
public:
OffloadEntryInfoDeviceGlobalVar()
: OffloadEntryInfo(OffloadingEntryInfoDeviceGlobalVar) {}
explicit OffloadEntryInfoDeviceGlobalVar(unsigned Order,
- OMPTargetGlobalVarEntryKind Flags)
- : OffloadEntryInfo(OffloadingEntryInfoDeviceGlobalVar, Order, Flags) {}
+ OMPTargetGlobalVarEntryKind Flags,
+ StringRef OrigName)
+ : OffloadEntryInfo(OffloadingEntryInfoDeviceGlobalVar, Order, Flags),
+ OrigName(OrigName) {}
explicit OffloadEntryInfoDeviceGlobalVar(
unsigned Order, llvm::Constant *Addr, CharUnits VarSize,
OMPTargetGlobalVarEntryKind Flags,
- llvm::GlobalValue::LinkageTypes Linkage)
+ llvm::GlobalValue::LinkageTypes Linkage, StringRef OrigName)
: OffloadEntryInfo(OffloadingEntryInfoDeviceGlobalVar, Order, Flags),
- VarSize(VarSize), Linkage(Linkage) {
+ VarSize(VarSize), Linkage(Linkage), OrigName(OrigName) {
setAddress(Addr);
}
@@ -684,16 +688,21 @@
static bool classof(const OffloadEntryInfo *Info) {
return Info->getKind() == OffloadingEntryInfoDeviceGlobalVar;
}
+ StringRef getOrigName() const { return OrigName; }
+ void setOrigName( StringRef Name ) { OrigName = Name; }
};
/// Initialize device global variable entry.
void initializeDeviceGlobalVarEntryInfo(StringRef Name,
OMPTargetGlobalVarEntryKind Flags,
- unsigned Order);
+ unsigned Order,
+ StringRef OrigName);
+ void enterDeviceGlobalVarMangledName(StringRef OrigName, StringRef Name);
/// Register device global variable entry.
void
- registerDeviceGlobalVarEntryInfo(StringRef VarName, llvm::Constant *Addr,
+ registerDeviceGlobalVarEntryInfo(StringRef VarName, StringRef OrigName,
+ llvm::Constant *Addr,
CharUnits VarSize,
OMPTargetGlobalVarEntryKind Flags,
llvm::GlobalValue::LinkageTypes Linkage);
@@ -707,7 +716,8 @@
OffloadDeviceGlobalVarEntryInfoActTy;
void actOnDeviceGlobalVarEntriesInfo(
const OffloadDeviceGlobalVarEntryInfoActTy &Action);
-
+ /// Return host mangled name
+ StringRef getOffloadEntryHostMangledName(StringRef VarName);
private:
// Storage for target region entries kind. The storage is to be indexed by
// file ID, device ID, parent function name and line number.
@@ -726,6 +736,8 @@
typedef llvm::StringMap<OffloadEntryInfoDeviceGlobalVar>
OffloadEntriesDeviceGlobalVarTy;
OffloadEntriesDeviceGlobalVarTy OffloadEntriesDeviceGlobalVar;
+ /// indexed by original name
+ llvm::StringMap<std::string> OffloadEntriesDeviceGlobalVarNameMap;
};
OffloadEntriesInfoManagerTy OffloadEntriesInfoManager;
@@ -1924,6 +1936,9 @@
/// Returns true if the variable is a local variable in untied task.
bool isLocalVarInUntiedTask(CodeGenFunction &CGF, const VarDecl *VD) const;
+
+ /// Returns the mangled name for declare target global
+ StringRef getHostMangledDeclareTargetGlobal( StringRef VarName );
};
/// Class supports emissionof SIMD-only code.
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -3032,19 +3032,35 @@
Action(D.first, F.first, P.first(), L.first, L.second);
}
+void CGOpenMPRuntime::OffloadEntriesInfoManagerTy::
+ enterDeviceGlobalVarMangledName(StringRef OrigName, StringRef Name) {
+ OffloadEntriesDeviceGlobalVarNameMap.try_emplace(OrigName, Name.str());
+}
+
void CGOpenMPRuntime::OffloadEntriesInfoManagerTy::
initializeDeviceGlobalVarEntryInfo(StringRef Name,
OMPTargetGlobalVarEntryKind Flags,
- unsigned Order) {
+ unsigned Order,
+ StringRef OrigName) {
assert(CGM.getLangOpts().OpenMPIsDevice && "Initialization of entries is "
"only required for the device "
"code generation.");
- OffloadEntriesDeviceGlobalVar.try_emplace(Name, Order, Flags);
+ OffloadEntriesDeviceGlobalVar.try_emplace(Name, Order, Flags, OrigName);
++OffloadingEntriesNum;
}
+StringRef CGOpenMPRuntime::OffloadEntriesInfoManagerTy::
+ getOffloadEntryHostMangledName(StringRef VarName) {
+ if ( OffloadEntriesDeviceGlobalVarNameMap.find(VarName) !=
+ OffloadEntriesDeviceGlobalVarNameMap.end() ) {
+ return OffloadEntriesDeviceGlobalVarNameMap[VarName];
+ }
+ return StringRef();
+}
+
void CGOpenMPRuntime::OffloadEntriesInfoManagerTy::
- registerDeviceGlobalVarEntryInfo(StringRef VarName, llvm::Constant *Addr,
+ registerDeviceGlobalVarEntryInfo(StringRef VarName, StringRef OrigName,
+ llvm::Constant *Addr,
CharUnits VarSize,
OMPTargetGlobalVarEntryKind Flags,
llvm::GlobalValue::LinkageTypes Linkage) {
@@ -3063,6 +3079,7 @@
Entry.setVarSize(VarSize);
Entry.setLinkage(Linkage);
Entry.setAddress(Addr);
+ Entry.setOrigName(OrigName);
} else {
if (hasDeviceGlobalVarEntryInfo(VarName)) {
auto &Entry = OffloadEntriesDeviceGlobalVar[VarName];
@@ -3075,7 +3092,7 @@
return;
}
OffloadEntriesDeviceGlobalVar.try_emplace(
- VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage);
+ VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage, OrigName);
++OffloadingEntriesNum;
}
}
@@ -3090,8 +3107,10 @@
void CGOpenMPRuntime::createOffloadEntry(
llvm::Constant *ID, llvm::Constant *Addr, uint64_t Size, int32_t Flags,
- llvm::GlobalValue::LinkageTypes Linkage) {
- OMPBuilder.emitOffloadingEntry(ID, Addr->getName(), Size, Flags);
+ llvm::GlobalValue::LinkageTypes Linkage,
+ StringRef MangledName) {
+ StringRef VarName = (MangledName.empty()) ? Addr->getName() : MangledName;
+ OMPBuilder.emitOffloadingEntry(ID, VarName, Size, Flags);
}
void CGOpenMPRuntime::createOffloadEntriesAndInfoMetadata() {
@@ -3184,10 +3203,11 @@
// - Entry 1 -> Mangled name of the variable.
// - Entry 2 -> Declare target kind.
// - Entry 3 -> Order the entry was created.
+ // - Entry 4 -> Original name of the variable.
// The first element of the metadata node is the kind.
llvm::Metadata *Ops[] = {
GetMDInt(E.getKind()), GetMDString(MangledName),
- GetMDInt(E.getFlags()), GetMDInt(E.getOrder())};
+ GetMDInt(E.getFlags()), GetMDInt(E.getOrder()), GetMDString(E.getOrigName())};
// Save this entry in the right position of the ordered entries array.
OrderedEntries[E.getOrder()] =
@@ -3218,7 +3238,8 @@
continue;
}
createOffloadEntry(CE->getID(), CE->getAddress(), /*Size=*/0,
- CE->getFlags(), llvm::GlobalValue::WeakAnyLinkage);
+ CE->getFlags(), llvm::GlobalValue::WeakAnyLinkage,
+ /*MangledName*/StringRef());
} else if (const auto *CE = dyn_cast<OffloadEntriesInfoManagerTy::
OffloadEntryInfoDeviceGlobalVar>(
std::get<0>(E))) {
@@ -3263,12 +3284,13 @@
// Hidden or internal symbols on the device are not externally visible. We
// should not attempt to register them by creating an offloading entry.
if (auto *GV = dyn_cast<llvm::GlobalValue>(CE->getAddress()))
- if (GV->hasLocalLinkage() || GV->hasHiddenVisibility())
+ if (GV->hasHiddenVisibility())
continue;
+ StringRef MangledName = std::get<2>(E);
createOffloadEntry(CE->getAddress(), CE->getAddress(),
CE->getVarSize().getQuantity(), Flags,
- CE->getLinkage());
+ CE->getLinkage(), MangledName);
} else {
llvm_unreachable("Unsupported entry kind.");
}
@@ -3338,12 +3360,20 @@
/*MangledName=*/GetMDString(1),
static_cast<OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryKind>(
/*Flags=*/GetMDInt(2)),
- /*Order=*/GetMDInt(3));
+ /*Order=*/GetMDInt(3),
+ /*OrigName=*/GetMDString(4));
+ OffloadEntriesInfoManager.enterDeviceGlobalVarMangledName(
+ /*OrigName=*/GetMDString(4),
+ /*MangledName=*/GetMDString(1));
break;
}
}
}
+StringRef CGOpenMPRuntime::getHostMangledDeclareTargetGlobal( StringRef VarName ) {
+ return OffloadEntriesInfoManager.getOffloadEntryHostMangledName(VarName);
+}
+
void CGOpenMPRuntime::emitKmpRoutineEntryT(QualType KmpInt32Ty) {
if (!KmpRoutineEntryPtrTy) {
// Build typedef kmp_int32 (* kmp_routine_entry_t)(kmp_int32, void *); type.
@@ -10750,11 +10780,29 @@
StringRef VarName;
CharUnits VarSize;
llvm::GlobalValue::LinkageTypes Linkage;
+ StringRef OrigName = VD->getName();
+ SmallString<256> Buffer;
+ llvm::raw_svector_ostream Out(Buffer);
if (*Res == OMPDeclareTargetDeclAttr::MT_To &&
!HasRequiresUnifiedSharedMemory) {
Flags = OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryTo;
- VarName = CGM.getMangledName(VD);
+ // We don't need to mangle the host side of declare target global variables but we
+ // need to create offload entry that matches the device side which gets mangled.
+ if (!CGM.getLangOpts().OpenMPIsDevice &&
+ !VD->isExternallyVisible()) {
+ VarName = OffloadEntriesInfoManager.getOffloadEntryHostMangledName(OrigName);
+ if ( VarName.empty() ) {
+ Out<<OrigName;
+ CGM.printPostfixForExternalizedDecl(Out, VD);
+ OffloadEntriesInfoManager.enterDeviceGlobalVarMangledName(OrigName, Buffer);
+ VarName = Buffer;
+ }
+ }
+ else {
+ VarName = CGM.getMangledName(VD);
+ }
+
if (VD->hasDefinition(CGM.getContext()) != VarDecl::DeclarationOnly) {
VarSize = CGM.getContext().getTypeSizeInChars(VD->getType());
assert(!VarSize.isZero() && "Expected non-zero size of the variable");
@@ -10801,7 +10849,7 @@
}
OffloadEntriesInfoManager.registerDeviceGlobalVarEntryInfo(
- VarName, Addr, VarSize, Flags, Linkage);
+ VarName, OrigName, Addr, VarSize, Flags, Linkage);
}
bool CGOpenMPRuntime::emitTargetGlobal(GlobalDecl GD) {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
