ssquare08 updated this revision to Diff 447901.
ssquare08 added a comment.
Herald added a project: OpenMP.
Herald added a subscriber: openmp-commits.
Adding a test and fixing
This adds a new runtime test and also address some comments.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D129694/new/
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_only_one_side_compilation.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
openmp/libomptarget/test/mapping/declare_target_static_var.c
Index: openmp/libomptarget/test/mapping/declare_target_static_var.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/mapping/declare_target_static_var.c
@@ -0,0 +1,21 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <stdio.h>
+
+#pragma omp declare target
+static int y;
+#pragma omp end declare target
+
+int main(void) {
+ y = 2;
+#pragma omp target update to(y)
+
+#pragma omp target
+ { y += 3; }
+
+#pragma omp target update from(y)
+
+ // CHECK: Declare target var update successful
+ printf("Declare target var update %s\n", (y == 5) ? "successful" : "failed");
+ return 0;
+}
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* @_ZN2ns1aE1, 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]+_l[0-9]+}}
+// HOST: @.omp_offloading.entry.x__static__{{[0-9a-z]+_[0-9a-z]+_l[0-9]+}}
// HOST-NOT: @.omp_offloading.entry.y
// HOST: @.omp_offloading.entry.z
C() : x(0) {}
Index: clang/test/OpenMP/declare_target_only_one_side_compilation.cpp
===================================================================
--- clang/test/OpenMP/declare_target_only_one_side_compilation.cpp
+++ clang/test/OpenMP/declare_target_only_one_side_compilation.cpp
@@ -58,7 +58,7 @@
// TODO: It is odd, probably wrong, that we don't mangle all variables.
// DEVICE-DAG: @G1 = {{.*}}global i32 0, align 4
-// DEVICE-DAG: @_ZL2G2 = internal {{.*}}global i32 0, align 4
+// DEVICE-DAG: @_ZL2G2 = {{.*}}global i32 0, align 4
// DEVICE-DAG: @G3 = {{.*}}global i32 0, align 4
// DEVICE-DAG: @_ZL2G4 = internal {{.*}}global i32 0, align 4
// DEVICE-DAG: @G5 = {{.*}}global i32 0, align 4
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__.+]] = 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
@@ -12959,20 +12959,19 @@
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
@@ -7294,6 +7294,7 @@
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
if (GV->isDeclaration())
return;
+
const VarDecl *VD = dyn_cast_or_null<VarDecl>(D);
if (VD) {
if (M.getLangOpts().CUDA) {
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -1519,6 +1519,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
@@ -4274,13 +4291,13 @@
// Handle things which are present even on external declarations.
if (D) {
- if (LangOpts.OpenMP && !LangOpts.OpenMPSimd)
- getOpenMPRuntime().registerTargetGlobalVariable(D, GV);
-
// FIXME: This code is overly simple and should be merged with other global
// handling.
GV->setConstant(isTypeConstant(D->getType(), false));
+ if (LangOpts.OpenMP && !LangOpts.OpenMPSimd)
+ getOpenMPRuntime().registerTargetGlobalVariable(D, GV);
+
GV->setAlignment(getContext().getDeclAlign(D).getAsAlign());
setLinkageForGV(GV, D);
@@ -4862,7 +4879,20 @@
!D->hasAttr<ConstInitAttr>())
Linkage = llvm::GlobalValue::InternalLinkage;
- GV->setLinkage(Linkage);
+ // Make sure any variable with OpenMP declare target is visible to the runtime
+ // except for constants and those with hidden visibility
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(D);
+ if (DevTy && (*DevTy == OMPDeclareTargetDeclAttr::DT_Any) &&
+ getLangOpts().OpenMPIsDevice && D && !GV->hasHiddenVisibility() &&
+ !GV->isConstant() &&
+ !getOpenMPRuntime().hasRequiresUnifiedSharedMemory()) {
+ GV->setLinkage(llvm::GlobalValue::ExternalLinkage);
+ GV->setDSOLocal(false);
+ } else {
+ GV->setLinkage(Linkage);
+ }
+
if (D->hasAttr<DLLImportAttr>())
GV->setDLLStorageClass(llvm::GlobalVariable::DLLImportStorageClass);
else if (D->hasAttr<DLLExportAttr>())
@@ -6968,6 +6998,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
@@ -1120,9 +1120,10 @@
}
void CGOpenMPRuntimeGPU::createOffloadEntry(llvm::Constant *ID,
- llvm::Constant *Addr,
- uint64_t Size, int32_t,
- llvm::GlobalValue::LinkageTypes) {
+ llvm::Constant *Addr, uint64_t Size,
+ int32_t,
+ 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,21 +662,24 @@
/// 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);
- }
+ }
CharUnits getVarSize() const { return VarSize; }
void setVarSize(CharUnits Size) { VarSize = Size; }
@@ -684,17 +688,20 @@
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,
- CharUnits VarSize,
+ registerDeviceGlobalVarEntryInfo(StringRef VarName, StringRef OrigName,
+ llvm::Constant *Addr, CharUnits VarSize,
OMPTargetGlobalVarEntryKind Flags,
llvm::GlobalValue::LinkageTypes Linkage);
/// Checks if the variable with the given name has been registered already.
@@ -707,6 +714,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
@@ -726,6 +735,8 @@
typedef llvm::StringMap<OffloadEntryInfoDeviceGlobalVar>
OffloadEntriesDeviceGlobalVarTy;
OffloadEntriesDeviceGlobalVarTy OffloadEntriesDeviceGlobalVar;
+ /// indexed by original name
+ llvm::StringMap<std::string> OffloadEntriesDeviceGlobalVarNameMap;
};
OffloadEntriesInfoManagerTy OffloadEntriesInfoManager;
@@ -1924,6 +1935,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,20 +3032,38 @@
Action(D.first, F.first, P.first(), L.first, L.second);
}
+void CGOpenMPRuntime::OffloadEntriesInfoManagerTy::
+ enterDeviceGlobalVarMangledName(StringRef OrigName, StringRef MangledName) {
+ if (!OrigName.equals(MangledName)) {
+ OffloadEntriesDeviceGlobalVarNameMap.try_emplace(OrigName,
+ MangledName.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,
- CharUnits VarSize,
+ registerDeviceGlobalVarEntryInfo(StringRef VarName, StringRef OrigName,
+ llvm::Constant *Addr, CharUnits VarSize,
OMPTargetGlobalVarEntryKind Flags,
llvm::GlobalValue::LinkageTypes Linkage) {
if (CGM.getLangOpts().OpenMPIsDevice) {
@@ -3063,6 +3081,7 @@
Entry.setVarSize(VarSize);
Entry.setLinkage(Linkage);
Entry.setAddress(Addr);
+ Entry.setOrigName(OrigName);
} else {
if (hasDeviceGlobalVarEntryInfo(VarName)) {
auto &Entry = OffloadEntriesDeviceGlobalVar[VarName];
@@ -3075,7 +3094,7 @@
return;
}
OffloadEntriesDeviceGlobalVar.try_emplace(
- VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage);
+ VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage, OrigName);
++OffloadingEntriesNum;
}
}
@@ -3090,8 +3109,9 @@
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 +3204,12 @@
// - 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())};
+ llvm::Metadata *Ops[] = {GetMDInt(E.getKind()),
+ GetMDString(MangledName),
+ 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 +3240,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))) {
@@ -3260,15 +3283,18 @@
break;
}
- // Hidden or internal symbols on the device are not externally visible. We
- // should not attempt to register them by creating an offloading entry.
+ // Hidden symbols on the device are not externally visible and constants
+ // don't need to be modified. 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() ||
+ dyn_cast<llvm::GlobalVariable>(GV)->isConstant())
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 +3364,21 @@
/*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 +10785,39 @@
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.
+ auto *GV = dyn_cast<llvm::GlobalValue>(Addr);
+ if (!CGM.getLangOpts().OpenMPIsDevice && !VD->isExternallyVisible() &&
+ !GV->hasHiddenVisibility() &&
+ !dyn_cast<llvm::GlobalVariable>(GV)->isConstant()) {
+ VarName =
+ OffloadEntriesInfoManager.getOffloadEntryHostMangledName(OrigName);
+ if (VarName.empty()) {
+ unsigned DeviceID;
+ unsigned FileID;
+ unsigned Line;
+ SourceLocation Loc = VD->getCanonicalDecl()->getBeginLoc();
+ getTargetEntryUniqueInfo(CGM.getContext(), Loc, DeviceID, FileID, Line);
+ {
+ Out << VD->getName() << "__static__" << llvm::format("%x", DeviceID)
+ << llvm::format("_%x_", FileID) << "l" << Line;
+ }
+ 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 +10864,7 @@
}
OffloadEntriesInfoManager.registerDeviceGlobalVarEntryInfo(
- VarName, Addr, VarSize, Flags, Linkage);
+ VarName, OrigName, Addr, VarSize, Flags, Linkage);
}
bool CGOpenMPRuntime::emitTargetGlobal(GlobalDecl GD) {
@@ -11187,7 +11250,6 @@
isa<OMPTargetExitDataDirective>(D) ||
isa<OMPTargetUpdateDirective>(D)) &&
"Expecting either target enter, exit data, or update directives.");
-
CodeGenFunction::OMPTargetDataInfo InputInfo;
llvm::Value *MapTypesArray = nullptr;
llvm::Value *MapNamesArray = nullptr;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
