r339805 - [OPENMP] FIx processing of declare target variables.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Wed Aug 15 12:45:13 PDT 2018
Author: abataev
Date: Wed Aug 15 12:45:12 2018
New Revision: 339805
URL: http://llvm.org/viewvc/llvm-project?rev=339805&view=rev
Log:
[OPENMP] FIx processing of declare target variables.
The compiler may produce unexpected error messages/crashes when declare
target variables were used. Patch fixes problems with the declarations
marked as declare target to or link.
Modified:
cfe/trunk/lib/AST/ASTContext.cpp
cfe/trunk/lib/CodeGen/CGExpr.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/Serialization/ASTReaderDecl.cpp
cfe/trunk/lib/Serialization/ASTWriter.cpp
cfe/trunk/lib/Serialization/ASTWriterDecl.cpp
cfe/trunk/test/OpenMP/declare_target_codegen.cpp
cfe/trunk/test/OpenMP/declare_target_link_codegen.cpp
Modified: cfe/trunk/lib/AST/ASTContext.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTContext.cpp?rev=339805&r1=339804&r2=339805&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ASTContext.cpp (original)
+++ cfe/trunk/lib/AST/ASTContext.cpp Wed Aug 15 12:45:12 2018
@@ -9774,6 +9774,12 @@ bool ASTContext::DeclMustBeEmitted(const
const auto *VD = cast<VarDecl>(D);
assert(VD->isFileVarDecl() && "Expected file scoped var");
+ // If the decl is marked as `declare target to`, it should be emitted for the
+ // host and for the device.
+ if (LangOpts.OpenMP &&
+ OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
+ return true;
+
if (VD->isThisDeclarationADefinition() == VarDecl::DeclarationOnly &&
!isMSStaticDataMemberInlineDefinition(VD))
return false;
@@ -9805,11 +9811,6 @@ bool ASTContext::DeclMustBeEmitted(const
if (DeclMustBeEmitted(BindingVD))
return true;
- // If the decl is marked as `declare target`, it should be emitted.
- if (const llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
- OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
- return *Res != OMPDeclareTargetDeclAttr::MT_Link;
-
return false;
}
Modified: cfe/trunk/lib/CodeGen/CGExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGExpr.cpp?rev=339805&r1=339804&r2=339805&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGExpr.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGExpr.cpp Wed Aug 15 12:45:12 2018
@@ -2270,18 +2270,14 @@ static LValue EmitThreadPrivateVarDeclLV
static Address emitDeclTargetLinkVarDeclLValue(CodeGenFunction &CGF,
const VarDecl *VD, QualType T) {
- for (const auto *D : VD->redecls()) {
- if (!VD->hasAttrs())
- continue;
- if (const auto *Attr = D->getAttr<OMPDeclareTargetDeclAttr>())
- if (Attr->getMapType() == OMPDeclareTargetDeclAttr::MT_Link) {
- QualType PtrTy = CGF.getContext().getPointerType(VD->getType());
- Address Addr =
- CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
- return CGF.EmitLoadOfPointer(Addr, PtrTy->castAs<PointerType>());
- }
- }
- return Address::invalid();
+ llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
+ OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
+ if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_To)
+ return Address::invalid();
+ assert(*Res == OMPDeclareTargetDeclAttr::MT_Link && "Expected link clause");
+ QualType PtrTy = CGF.getContext().getPointerType(VD->getType());
+ Address Addr = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
+ return CGF.EmitLoadOfPointer(Addr, PtrTy->castAs<PointerType>());
}
Address
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=339805&r1=339804&r2=339805&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Wed Aug 15 12:45:12 2018
@@ -2622,7 +2622,7 @@ bool CGOpenMPRuntime::emitDeclareTargetV
Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link)
- return false;
+ return CGM.getLangOpts().OpenMPIsDevice;
VD = VD->getDefinition(CGM.getContext());
if (VD && !DeclareTargetWithDefinition.insert(VD).second)
return CGM.getLangOpts().OpenMPIsDevice;
@@ -8089,8 +8089,7 @@ bool CGOpenMPRuntime::emitTargetGlobalVa
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(
cast<VarDecl>(GD.getDecl()));
if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link) {
- if (CGM.getContext().DeclMustBeEmitted(GD.getDecl()))
- DeferredGlobalVariables.insert(cast<VarDecl>(GD.getDecl()));
+ DeferredGlobalVariables.insert(cast<VarDecl>(GD.getDecl()));
return true;
}
return false;
@@ -8154,10 +8153,14 @@ void CGOpenMPRuntime::emitDeferredTarget
for (const VarDecl *VD : DeferredGlobalVariables) {
llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
- if (Res) {
- assert(*Res != OMPDeclareTargetDeclAttr::MT_Link &&
- "Implicit declare target variables must be only to().");
+ if (!Res)
+ continue;
+ if (*Res == OMPDeclareTargetDeclAttr::MT_To) {
CGM.EmitGlobal(VD);
+ } else {
+ assert(*Res == OMPDeclareTargetDeclAttr::MT_Link &&
+ "Expected to or link clauses.");
+ (void)CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
}
}
}
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=339805&r1=339804&r2=339805&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Wed Aug 15 12:45:12 2018
@@ -2004,7 +2004,8 @@ bool CodeGenModule::MayBeEmittedEagerly(
// codegen for global variables, because they may be marked as threadprivate.
if (LangOpts.OpenMP && LangOpts.OpenMPUseTLS &&
getContext().getTargetInfo().isTLSSupported() && isa<VarDecl>(Global) &&
- !isTypeConstant(Global->getType(), false))
+ !isTypeConstant(Global->getType(), false) &&
+ !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(Global))
return false;
return true;
@@ -2155,6 +2156,20 @@ void CodeGenModule::EmitGlobal(GlobalDec
if (!MustEmitForCuda &&
VD->isThisDeclarationADefinition() != VarDecl::Definition &&
!Context.isMSStaticDataMemberInlineDefinition(VD)) {
+ if (LangOpts.OpenMP) {
+ // Emit declaration of the must-be-emitted declare target variable.
+ if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
+ OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
+ if (*Res == OMPDeclareTargetDeclAttr::MT_To) {
+ (void)GetAddrOfGlobalVar(VD);
+ } else {
+ assert(*Res == OMPDeclareTargetDeclAttr::MT_Link &&
+ "link claue expected.");
+ (void)getOpenMPRuntime().getAddrOfDeclareTargetLink(VD);
+ }
+ return;
+ }
+ }
// If this declaration may have caused an inline variable definition to
// change linkage, make sure that it's emitted.
if (Context.getInlineVariableDefinitionKind(VD) ==
Modified: cfe/trunk/lib/Serialization/ASTReaderDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTReaderDecl.cpp?rev=339805&r1=339804&r2=339805&view=diff
==============================================================================
--- cfe/trunk/lib/Serialization/ASTReaderDecl.cpp (original)
+++ cfe/trunk/lib/Serialization/ASTReaderDecl.cpp Wed Aug 15 12:45:12 2018
@@ -2708,7 +2708,8 @@ static bool isConsumerInterestedIn(ASTCo
return !D->getDeclContext()->isFunctionOrMethod();
if (const auto *Var = dyn_cast<VarDecl>(D))
return Var->isFileVarDecl() &&
- Var->isThisDeclarationADefinition() == VarDecl::Definition;
+ (Var->isThisDeclarationADefinition() == VarDecl::Definition ||
+ OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(Var));
if (const auto *Func = dyn_cast<FunctionDecl>(D))
return Func->doesThisDeclarationHaveABody() || HasBody;
@@ -4385,6 +4386,12 @@ void ASTDeclReader::UpdateDecl(Decl *D,
}
case UPD_DECL_MARKED_OPENMP_DECLARETARGET:
+ D->addAttr(OMPDeclareTargetDeclAttr::CreateImplicit(
+ Reader.getContext(),
+ static_cast<OMPDeclareTargetDeclAttr::MapTypeTy>(Record.readInt()),
+ ReadSourceRange()));
+ break;
+
case UPD_ADDED_ATTR_TO_RECORD:
AttrVec Attrs;
Record.readAttributes(Attrs);
Modified: cfe/trunk/lib/Serialization/ASTWriter.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTWriter.cpp?rev=339805&r1=339804&r2=339805&view=diff
==============================================================================
--- cfe/trunk/lib/Serialization/ASTWriter.cpp (original)
+++ cfe/trunk/lib/Serialization/ASTWriter.cpp Wed Aug 15 12:45:12 2018
@@ -5296,6 +5296,7 @@ void ASTWriter::WriteDeclUpdatesBlocks(R
break;
case UPD_DECL_MARKED_OPENMP_DECLARETARGET:
+ Record.push_back(D->getAttr<OMPDeclareTargetDeclAttr>()->getMapType());
Record.AddSourceRange(
D->getAttr<OMPDeclareTargetDeclAttr>()->getRange());
break;
Modified: cfe/trunk/lib/Serialization/ASTWriterDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTWriterDecl.cpp?rev=339805&r1=339804&r2=339805&view=diff
==============================================================================
--- cfe/trunk/lib/Serialization/ASTWriterDecl.cpp (original)
+++ cfe/trunk/lib/Serialization/ASTWriterDecl.cpp Wed Aug 15 12:45:12 2018
@@ -2237,8 +2237,7 @@ static bool isRequiredDecl(const Decl *D
// File scoped assembly or obj-c or OMP declare target implementation must be
// seen.
- if (isa<FileScopeAsmDecl>(D) || isa<ObjCImplDecl>(D) ||
- D->hasAttr<OMPDeclareTargetDeclAttr>())
+ if (isa<FileScopeAsmDecl>(D) || isa<ObjCImplDecl>(D))
return true;
if (WritingModule && (isa<VarDecl>(D) || isa<ImportDecl>(D))) {
Modified: cfe/trunk/test/OpenMP/declare_target_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_codegen.cpp?rev=339805&r1=339804&r2=339805&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/declare_target_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/declare_target_codegen.cpp Wed Aug 15 12:45:12 2018
@@ -13,6 +13,15 @@
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
// CHECK-NOT: define {{.*}}{{baz1|baz4|maini1}}
+// CHECK-NOT: @{{hhh|ggg|fff|eee}} =
+// CHECK-DAG: @aaa = external global i32,
+// CHECK-DAG: @bbb = global i32 0,
+// CHECK-DAG: @ccc = external global i32,
+// CHECK-DAG: @ddd = global i32 0,
+// CHECK-DAG: @hhh_decl_tgt_link_ptr = common global i32* null
+// CHECK-DAG: @ggg_decl_tgt_link_ptr = common global i32* null
+// CHECK-DAG: @fff_decl_tgt_link_ptr = common global i32* null
+// CHECK-DAG: @eee_decl_tgt_link_ptr = common global i32* null
// CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,
// CHECK-DAG: @b = global i32 15,
// CHECK-DAG: @d = global i32 0,
@@ -21,17 +30,30 @@
// CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
// CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
// CHECK-DAG: @out_decl_target = global i32 0,
-// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+56]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+57]]_ctor to i8*)],
+// CHECK-DAG: @llvm.used = appending global [6 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+69]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+70]]_ctor to i8*),
// CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)],
// CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}()
// CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}})
// CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}})
-// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+50]]_ctor()
+// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+63]]_ctor()
#ifndef HEADER
#define HEADER
+#pragma omp declare target
+extern int aaa;
+int bbb = 0;
+extern int ccc;
+int ddd = 0;
+#pragma omp end declare target
+
+extern int eee;
+int fff = 0;
+extern int ggg;
+int hhh = 0;
+#pragma omp declare target link(eee, fff, ggg, hhh)
+
int out_decl_target = 0;
#pragma omp declare target
void lambda () {
@@ -86,7 +108,7 @@ int bar() { return 1 + foo() + bar() + b
int maini1() {
int a;
- static long aa = 32;
+ static long aa = 32 + bbb + ccc + fff + ggg;
// CHECK-DAG: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* dereferenceable{{.*}}, i64 {{.*}}, i64 {{.*}})
#pragma omp target map(tofrom \
: a, b)
Modified: cfe/trunk/test/OpenMP/declare_target_link_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_link_codegen.cpp?rev=339805&r1=339804&r2=339805&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/declare_target_link_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/declare_target_link_codegen.cpp Wed Aug 15 12:45:12 2018
@@ -17,10 +17,10 @@
#ifndef HEADER
#define HEADER
-// HOST: @c = external global i32,
+// HOST-DAG: @c = external global i32,
+// HOST-DAG: @c_decl_tgt_link_ptr = global i32* @c
// DEVICE-NOT: @c =
// DEVICE: @c_decl_tgt_link_ptr = common global i32* null
-// HOST: @c_decl_tgt_link_ptr = global i32* @c
// HOST: [[SIZES:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 4]
// HOST: [[MAPTYPES:@.+]] = private unnamed_addr constant [2 x i64] [i64 35, i64 531]
// HOST: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"c_decl_tgt_link_ptr\00"
More information about the cfe-commits
mailing list