From d01b74974ba2e87cfc1f05a08ccad54feb9ceb49 Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Wed, 15 Aug 2018 19:45:12 +0000 Subject: [PATCH] [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. llvm-svn: 339805 --- clang/lib/AST/ASTContext.cpp | 11 +++++---- clang/lib/CodeGen/CGExpr.cpp | 20 +++++++--------- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 15 +++++++----- clang/lib/CodeGen/CodeGenModule.cpp | 17 +++++++++++++- clang/lib/Serialization/ASTReaderDecl.cpp | 9 +++++++- clang/lib/Serialization/ASTWriter.cpp | 1 + clang/lib/Serialization/ASTWriterDecl.cpp | 3 +-- clang/test/OpenMP/declare_target_codegen.cpp | 28 ++++++++++++++++++++--- clang/test/OpenMP/declare_target_link_codegen.cpp | 4 ++-- 9 files changed, 76 insertions(+), 32 deletions(-) diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 5f94cf2..91e9f3e 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -9774,6 +9774,12 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { const auto *VD = cast(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 Decl *D) { if (DeclMustBeEmitted(BindingVD)) return true; - // If the decl is marked as `declare target`, it should be emitted. - if (const llvm::Optional Res = - OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) - return *Res != OMPDeclareTargetDeclAttr::MT_Link; - return false; } diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index 336fdd0..605c01d 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -2270,18 +2270,14 @@ static LValue EmitThreadPrivateVarDeclLValue( 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()) - 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()); - } - } - return Address::invalid(); + llvm::Optional 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()); } Address diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 245eacc..0072bdf 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -2622,7 +2622,7 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD, Optional 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::emitTargetGlobalVariable(GlobalDecl GD) { OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration( cast(GD.getDecl())); if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link) { - if (CGM.getContext().DeclMustBeEmitted(GD.getDecl())) - DeferredGlobalVariables.insert(cast(GD.getDecl())); + DeferredGlobalVariables.insert(cast(GD.getDecl())); return true; } return false; @@ -8154,10 +8153,14 @@ void CGOpenMPRuntime::emitDeferredTargetDecls() const { for (const VarDecl *VD : DeferredGlobalVariables) { llvm::Optional 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); } } } diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 72fa799..76527b7 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2004,7 +2004,8 @@ bool CodeGenModule::MayBeEmittedEagerly(const ValueDecl *Global) { // codegen for global variables, because they may be marked as threadprivate. if (LangOpts.OpenMP && LangOpts.OpenMPUseTLS && getContext().getTargetInfo().isTLSSupported() && isa(Global) && - !isTypeConstant(Global->getType(), false)) + !isTypeConstant(Global->getType(), false) && + !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(Global)) return false; return true; @@ -2155,6 +2156,20 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) { 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 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) == diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp index 2b8cfac..47484f8 100644 --- a/clang/lib/Serialization/ASTReaderDecl.cpp +++ b/clang/lib/Serialization/ASTReaderDecl.cpp @@ -2708,7 +2708,8 @@ static bool isConsumerInterestedIn(ASTContext &Ctx, Decl *D, bool HasBody) { return !D->getDeclContext()->isFunctionOrMethod(); if (const auto *Var = dyn_cast(D)) return Var->isFileVarDecl() && - Var->isThisDeclarationADefinition() == VarDecl::Definition; + (Var->isThisDeclarationADefinition() == VarDecl::Definition || + OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(Var)); if (const auto *Func = dyn_cast(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(Record.readInt()), + ReadSourceRange())); + break; + case UPD_ADDED_ATTR_TO_RECORD: AttrVec Attrs; Record.readAttributes(Attrs); diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index ea3a75d..99606fc 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -5296,6 +5296,7 @@ void ASTWriter::WriteDeclUpdatesBlocks(RecordDataImpl &OffsetsRecord) { break; case UPD_DECL_MARKED_OPENMP_DECLARETARGET: + Record.push_back(D->getAttr()->getMapType()); Record.AddSourceRange( D->getAttr()->getRange()); break; diff --git a/clang/lib/Serialization/ASTWriterDecl.cpp b/clang/lib/Serialization/ASTWriterDecl.cpp index 2d90c98..8eb6c61 100644 --- a/clang/lib/Serialization/ASTWriterDecl.cpp +++ b/clang/lib/Serialization/ASTWriterDecl.cpp @@ -2237,8 +2237,7 @@ static bool isRequiredDecl(const Decl *D, ASTContext &Context, // File scoped assembly or obj-c or OMP declare target implementation must be // seen. - if (isa(D) || isa(D) || - D->hasAttr()) + if (isa(D) || isa(D)) return true; if (WritingModule && (isa(D) || isa(D))) { diff --git a/clang/test/OpenMP/declare_target_codegen.cpp b/clang/test/OpenMP/declare_target_codegen.cpp index fbeb6fb..f4ef422 100644 --- a/clang/test/OpenMP/declare_target_codegen.cpp +++ b/clang/test/OpenMP/declare_target_codegen.cpp @@ -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() + baz1() + baz2(); } 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) diff --git a/clang/test/OpenMP/declare_target_link_codegen.cpp b/clang/test/OpenMP/declare_target_link_codegen.cpp index fdecd72..564581c 100644 --- a/clang/test/OpenMP/declare_target_link_codegen.cpp +++ b/clang/test/OpenMP/declare_target_link_codegen.cpp @@ -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" -- 2.7.4