gtbercea updated this revision to Diff 205667.
gtbercea marked an inline comment as done.
gtbercea added a comment.
- Merge MT_Link and MT_To with unified memory cases.
- Transform switch into if statements.
- Fix declare target attribute checks.
Repository:
rC Clang
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D63108/new/
https://reviews.llvm.org/D63108
Files:
lib/CodeGen/CGDeclCXX.cpp
lib/CodeGen/CGExpr.cpp
lib/CodeGen/CGOpenMPRuntime.cpp
lib/CodeGen/CGOpenMPRuntime.h
lib/CodeGen/CodeGenModule.cpp
test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp
Index: test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp
===================================================================
--- test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp
+++ test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp
@@ -8,16 +8,18 @@
#define N 1000
double var = 10.0;
+double to_var = 20.0;
#pragma omp requires unified_shared_memory
#pragma omp declare target link(var)
+#pragma omp declare target to(to_var)
int bar(int n){
double sum = 0;
#pragma omp target
for(int i = 0; i < n; i++) {
- sum += var;
+ sum += var + to_var;
}
return sum;
@@ -26,9 +28,20 @@
// CHECK: [[VAR:@.+]] = global double 1.000000e+01
// CHECK: [[VAR_DECL_TGT_LINK_PTR:@.+]] = global double* [[VAR]]
+// CHECK: [[TO_VAR:@.+]] = global double 2.000000e+01
+// CHECK: [[VAR_DECL_TGT_TO_PTR:@.+]] = global double* [[TO_VAR]]
+
// CHECK: [[OFFLOAD_SIZES:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 8]
// CHECK: [[OFFLOAD_MAPTYPES:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
+// CHECK: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [22 x i8]
+// CHECK: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (double** [[VAR_DECL_TGT_LINK_PTR]] to i8*), i8* getelementptr inbounds ([22 x i8], [22 x i8]* [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME]], i32 0, i32 0), i64 8, i32 1, i32 0 }, section ".omp_offloading.entries"
+
+// CHECK: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [23 x i8]
+// CHECK: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (double** [[VAR_DECL_TGT_TO_PTR]] to i8*), i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME]], i32 0, i32 0), i64 8, i32 0, i32 0 }, section ".omp_offloading.entries"
+
+// CHECK: @llvm.used = appending global [2 x i8*] [i8* bitcast (double** [[VAR_DECL_TGT_LINK_PTR]] to i8*), i8* bitcast (double** [[VAR_DECL_TGT_TO_PTR]] to i8*)], section "llvm.metadata"
+
// CHECK: [[N_CASTED:%.+]] = alloca i64
// CHECK: [[SUM_CASTED:%.+]] = alloca i64
Index: lib/CodeGen/CodeGenModule.cpp
===================================================================
--- lib/CodeGen/CodeGenModule.cpp
+++ lib/CodeGen/CodeGenModule.cpp
@@ -2475,13 +2475,18 @@
// Emit declaration of the must-be-emitted declare target variable.
if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
- if (*Res == OMPDeclareTargetDeclAttr::MT_To) {
+ bool UnifiedMemoryEnabled =
+ getOpenMPRuntime().hasRequiresUnifiedSharedMemory();
+ if (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ !UnifiedMemoryEnabled)
(void)GetAddrOfGlobalVar(VD);
- } else {
- assert(*Res == OMPDeclareTargetDeclAttr::MT_Link &&
- "link claue expected.");
- (void)getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
- }
+ else if (*Res == OMPDeclareTargetDeclAttr::MT_Link ||
+ (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ UnifiedMemoryEnabled))
+ (void)getOpenMPRuntime().getAddrOfDeclareTargetClause(VD);
+ else
+ llvm_unreachable("Link or to clause expected!");
+
return;
}
}
Index: lib/CodeGen/CGOpenMPRuntime.h
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.h
+++ lib/CodeGen/CGOpenMPRuntime.h
@@ -1121,8 +1121,8 @@
SourceLocation Loc);
/// Returns the address of the variable marked as declare target with link
- /// clause.
- virtual Address getAddrOfDeclareTargetLink(const VarDecl *VD);
+ /// clause OR as declare target with to clause and unified memory.
+ virtual Address getAddrOfDeclareTargetClause(const VarDecl *VD);
/// Emit a code for initialization of threadprivate variable. It emits
/// a call to runtime library which adds initial value to the newly created
Index: lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.cpp
+++ lib/CodeGen/CGOpenMPRuntime.cpp
@@ -2552,16 +2552,21 @@
return CGM.CreateRuntimeFunction(FnTy, Name);
}
-Address CGOpenMPRuntime::getAddrOfDeclareTargetLink(const VarDecl *VD) {
+Address CGOpenMPRuntime::getAddrOfDeclareTargetClause(const VarDecl *VD) {
if (CGM.getLangOpts().OpenMPSimd)
return Address::invalid();
llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
- if (Res && *Res == OMPDeclareTargetDeclAttr::MT_Link) {
+ if (Res && (*Res == OMPDeclareTargetDeclAttr::MT_Link ||
+ (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ HasRequiresUnifiedSharedMemory))) {
SmallString<64> PtrName;
{
llvm::raw_svector_ostream OS(PtrName);
- OS << CGM.getMangledName(GlobalDecl(VD)) << "_decl_tgt_link_ptr";
+ if (*Res == OMPDeclareTargetDeclAttr::MT_Link)
+ OS << CGM.getMangledName(GlobalDecl(VD)) << "_decl_tgt_link_ptr";
+ else
+ OS << CGM.getMangledName(GlobalDecl(VD)) << "_decl_tgt_to_ptr";
}
llvm::Value *Ptr = CGM.getModule().getNamedValue(PtrName);
if (!Ptr) {
@@ -2778,7 +2783,9 @@
bool PerformInit) {
Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
- if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link)
+ if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link ||
+ (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ HasRequiresUnifiedSharedMemory))
return CGM.getLangOpts().OpenMPIsDevice;
VD = VD->getDefinition(CGM.getContext());
if (VD && !DeclareTargetWithDefinition.insert(CGM.getMangledName(VD)).second)
@@ -4194,6 +4201,9 @@
CE->getFlags());
switch (Flags) {
case OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryTo: {
+ if (CGM.getLangOpts().OpenMPIsDevice &&
+ CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())
+ continue;
if (!CE->getAddress()) {
unsigned DiagID = CGM.getDiags().getCustomDiagID(
DiagnosticsEngine::Error,
@@ -7452,7 +7462,7 @@
// Track if the map information being generated is the first for a capture.
bool IsCaptureFirstInfo = IsFirstComponentList;
- bool IsLink = false; // Is this variable a "declare target link"?
+ bool IsLinkOrToClause = false; // Is this variable a "declare target link"?
// Scan the components from the base to the complete expression.
auto CI = Components.rbegin();
@@ -7482,11 +7492,14 @@
if (const auto *VD =
dyn_cast_or_null<VarDecl>(I->getAssociatedDeclaration())) {
if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
- OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
- if (*Res == OMPDeclareTargetDeclAttr::MT_Link) {
- IsLink = true;
- BP = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
+ OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
+ if ((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
+ (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())) {
+ IsLinkOrToClause = true;
+ BP = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetClause(VD);
}
+ }
}
// If the variable is a pointer and is being dereferenced (i.e. is not
@@ -7652,7 +7665,8 @@
// (there is a set of entries for each capture).
OpenMPOffloadMappingFlags Flags = getMapTypeBits(
MapType, MapModifiers, IsImplicit,
- !IsExpressionFirstInfo || IsLink, IsCaptureFirstInfo && !IsLink);
+ !IsExpressionFirstInfo || IsLinkOrToClause,
+ IsCaptureFirstInfo && !IsLinkOrToClause);
if (!IsExpressionFirstInfo) {
// If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
@@ -9124,7 +9138,9 @@
llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(
cast<VarDecl>(GD.getDecl()));
- if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link) {
+ if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link ||
+ (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ HasRequiresUnifiedSharedMemory)) {
DeferredGlobalVariables.insert(cast<VarDecl>(GD.getDecl()));
return true;
}
@@ -9183,8 +9199,9 @@
StringRef VarName;
CharUnits VarSize;
llvm::GlobalValue::LinkageTypes Linkage;
- switch (*Res) {
- case OMPDeclareTargetDeclAttr::MT_To:
+
+ if (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ !HasRequiresUnifiedSharedMemory) {
Flags = OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryTo;
VarName = CGM.getMangledName(VD);
if (VD->hasDefinition(CGM.getContext()) != VarDecl::DeclarationOnly) {
@@ -9207,20 +9224,26 @@
CGM.addCompilerUsedGlobal(GVAddrRef);
}
}
- break;
- case OMPDeclareTargetDeclAttr::MT_Link:
+ } else if ((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
+ (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ HasRequiresUnifiedSharedMemory)) {
Flags = OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryLink;
+ if (*Res == OMPDeclareTargetDeclAttr::MT_To)
+ Flags = OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryTo;
+
if (CGM.getLangOpts().OpenMPIsDevice) {
VarName = Addr->getName();
Addr = nullptr;
} else {
- VarName = getAddrOfDeclareTargetLink(VD).getName();
- Addr = cast<llvm::Constant>(getAddrOfDeclareTargetLink(VD).getPointer());
+ VarName = getAddrOfDeclareTargetClause(VD).getName();
+ Addr = cast<llvm::Constant>(getAddrOfDeclareTargetClause(VD).getPointer());
}
VarSize = CGM.getPointerSize();
Linkage = llvm::GlobalValue::WeakAnyLinkage;
- break;
+ } else {
+ llvm_unreachable("Declare target attribute must be to or link.");
}
+
OffloadEntriesInfoManager.registerDeviceGlobalVarEntryInfo(
VarName, Addr, VarSize, Flags, Linkage);
}
@@ -9239,12 +9262,14 @@
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
if (!Res)
continue;
- if (*Res == OMPDeclareTargetDeclAttr::MT_To) {
+ if (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ !HasRequiresUnifiedSharedMemory) {
CGM.EmitGlobal(VD);
} else {
- assert(*Res == OMPDeclareTargetDeclAttr::MT_Link &&
+ assert((*Res == OMPDeclareTargetDeclAttr::MT_Link ||
+ *Res == OMPDeclareTargetDeclAttr::MT_To) &&
"Expected to or link clauses.");
- (void)CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
+ (void)CGM.getOpenMPRuntime().getAddrOfDeclareTargetClause(VD);
}
}
}
Index: lib/CodeGen/CGExpr.cpp
===================================================================
--- lib/CodeGen/CGExpr.cpp
+++ lib/CodeGen/CGExpr.cpp
@@ -2295,15 +2295,22 @@
return CGF.MakeAddrLValue(Addr, T, AlignmentSource::Decl);
}
-static Address emitDeclTargetLinkVarDeclLValue(CodeGenFunction &CGF,
- const VarDecl *VD, QualType T) {
+static Address emitDeclTargetVarDeclLValue(CodeGenFunction &CGF,
+ const VarDecl *VD, QualType T) {
llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
- if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_To)
+ // Return an invalid address if variable is MT_To and unified
+ // memory is not enabled. For all other cases: MT_Link and
+ // MT_To with unified memory, return a valid address.
+ if (!Res || (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ !CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory()))
return Address::invalid();
- assert(*Res == OMPDeclareTargetDeclAttr::MT_Link && "Expected link clause");
+ assert(((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
+ (*Res == OMPDeclareTargetDeclAttr::MT_To &&
+ CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())) &&
+ "Expected link clause OR to clause with unified memory enabled.");
QualType PtrTy = CGF.getContext().getPointerType(VD->getType());
- Address Addr = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
+ Address Addr = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetClause(VD);
return CGF.EmitLoadOfPointer(Addr, PtrTy->castAs<PointerType>());
}
@@ -2359,7 +2366,7 @@
// Check if the variable is marked as declare target with link clause in
// device codegen.
if (CGF.getLangOpts().OpenMPIsDevice) {
- Address Addr = emitDeclTargetLinkVarDeclLValue(CGF, VD, T);
+ Address Addr = emitDeclTargetVarDeclLValue(CGF, VD, T);
if (Addr.isValid())
return CGF.MakeAddrLValue(Addr, T, AlignmentSource::Decl);
}
Index: lib/CodeGen/CGDeclCXX.cpp
===================================================================
--- lib/CodeGen/CGDeclCXX.cpp
+++ lib/CodeGen/CGDeclCXX.cpp
@@ -74,7 +74,7 @@
// bails even if the attribute is not present.
if (D.isNoDestroy(CGF.getContext()))
return;
-
+
CodeGenModule &CGM = CGF.CGM;
// FIXME: __attribute__((cleanup)) ?
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
